From 79ab76e156d4bde937335be784cdf220294600d5 Mon Sep 17 00:00:00 2001
From: Brecht Van Lommel <brecht@blender.org>
Date: Thu, 28 Jul 2022 19:57:30 +0200
Subject: Cleanup: simplifications and consistency for vector types

* OneAPI: remove separate float3 definition
* OneAPI: disable operator[] to match other GPUs
* OneAPI: make int3 compact to match other GPUs
* Use #pragma once
* Add __KERNEL_NATIVE_VECTOR_TYPES__ to simplify checks
* Remove unused vector3
---
 intern/cycles/kernel/CMakeLists.txt         |  2 --
 intern/cycles/kernel/device/cpu/compat.h    | 34 -----------------------------
 intern/cycles/kernel/device/oneapi/compat.h | 16 ++------------
 intern/cycles/util/CMakeLists.txt           |  2 --
 intern/cycles/util/types.h                  | 11 ++++++----
 intern/cycles/util/types_float2.h           | 11 +++++-----
 intern/cycles/util/types_float2_impl.h      | 15 +++++--------
 intern/cycles/util/types_float3.h           | 28 +++++++++++++++---------
 intern/cycles/util/types_float3_impl.h      | 31 ++++++++++++++------------
 intern/cycles/util/types_float4.h           | 11 +++++-----
 intern/cycles/util/types_float4_impl.h      | 15 +++++--------
 intern/cycles/util/types_float8.h           |  7 ++----
 intern/cycles/util/types_float8_impl.h      |  9 +-------
 intern/cycles/util/types_int2.h             | 11 +++++-----
 intern/cycles/util/types_int2_impl.h        | 11 +++++-----
 intern/cycles/util/types_int3.h             | 23 +++++++++++--------
 intern/cycles/util/types_int3_impl.h        | 32 ++++++++++++++-------------
 intern/cycles/util/types_int4.h             | 11 +++++-----
 intern/cycles/util/types_int4_impl.h        | 15 +++++--------
 intern/cycles/util/types_uchar2.h           | 11 +++++-----
 intern/cycles/util/types_uchar2_impl.h      |  6 +++--
 intern/cycles/util/types_uchar3.h           |  6 +++--
 intern/cycles/util/types_uchar3_impl.h      |  6 +++--
 intern/cycles/util/types_uchar4.h           | 11 +++++-----
 intern/cycles/util/types_uchar4_impl.h      |  6 +++--
 intern/cycles/util/types_uint2.h            | 11 +++++-----
 intern/cycles/util/types_uint2_impl.h       | 11 +++++-----
 intern/cycles/util/types_uint3.h            | 11 +++++-----
 intern/cycles/util/types_uint3_impl.h       | 11 +++++-----
 intern/cycles/util/types_uint4.h            | 11 +++++-----
 intern/cycles/util/types_uint4_impl.h       | 11 +++++-----
 intern/cycles/util/types_ushort4.h          |  2 +-
 intern/cycles/util/types_vector3.h          | 26 ----------------------
 intern/cycles/util/types_vector3_impl.h     | 30 -------------------------
 34 files changed, 175 insertions(+), 290 deletions(-)
 delete mode 100644 intern/cycles/util/types_vector3.h
 delete mode 100644 intern/cycles/util/types_vector3_impl.h

(limited to 'intern')

diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 8ecdac6ee27..dfcd75a135e 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -355,8 +355,6 @@ set(SRC_UTIL_HEADERS
   ../util/types_uint4.h
   ../util/types_uint4_impl.h
   ../util/types_ushort4.h
-  ../util/types_vector3.h
-  ../util/types_vector3_impl.h
 )
 
 set(LIB
diff --git a/intern/cycles/kernel/device/cpu/compat.h b/intern/cycles/kernel/device/cpu/compat.h
index 631e55e0d42..1e3e790ca1f 100644
--- a/intern/cycles/kernel/device/cpu/compat.h
+++ b/intern/cycles/kernel/device/cpu/compat.h
@@ -33,38 +33,4 @@ CCL_NAMESPACE_BEGIN
 
 #define kernel_assert(cond) assert(cond)
 
-/* Macros to handle different memory storage on different devices */
-
-#ifdef __KERNEL_SSE2__
-typedef vector3<sseb> sse3b;
-typedef vector3<ssef> sse3f;
-typedef vector3<ssei> sse3i;
-
-ccl_device_inline void print_sse3b(const char *label, sse3b &a)
-{
-  print_sseb(label, a.x);
-  print_sseb(label, a.y);
-  print_sseb(label, a.z);
-}
-
-ccl_device_inline void print_sse3f(const char *label, sse3f &a)
-{
-  print_ssef(label, a.x);
-  print_ssef(label, a.y);
-  print_ssef(label, a.z);
-}
-
-ccl_device_inline void print_sse3i(const char *label, sse3i &a)
-{
-  print_ssei(label, a.x);
-  print_ssei(label, a.y);
-  print_ssei(label, a.z);
-}
-
-#  if defined(__KERNEL_AVX__) || defined(__KERNEL_AVX2__)
-typedef vector3<avxf> avx3f;
-#  endif
-
-#endif
-
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/device/oneapi/compat.h b/intern/cycles/kernel/device/oneapi/compat.h
index 1b25259bcf5..ccd141aa892 100644
--- a/intern/cycles/kernel/device/oneapi/compat.h
+++ b/intern/cycles/kernel/device/oneapi/compat.h
@@ -149,25 +149,13 @@ void oneapi_kernel_##name(KernelGlobalsGPU *ccl_restrict kg, \
 /* clang-format on */
 
 /* Types */
+
 /* It's not possible to use sycl types like sycl::float3, sycl::int3, etc
- * because these types have different interfaces from blender version */
+ * because these types have different interfaces from blender version. */
 
 using uchar = unsigned char;
 using sycl::half;
 
-struct float3 {
-  float x, y, z;
-};
-
-ccl_always_inline float3 make_float3(float x, float y, float z)
-{
-  return {x, y, z};
-}
-ccl_always_inline float3 make_float3(float x)
-{
-  return {x, x, x};
-}
-
 /* math functions */
 #define fabsf(x) sycl::fabs((x))
 #define copysignf(x, y) sycl::copysign((x), (y))
diff --git a/intern/cycles/util/CMakeLists.txt b/intern/cycles/util/CMakeLists.txt
index 9bc9f00e142..b33bad79e6c 100644
--- a/intern/cycles/util/CMakeLists.txt
+++ b/intern/cycles/util/CMakeLists.txt
@@ -129,8 +129,6 @@ set(SRC_HEADERS
   types_uint4.h
   types_uint4_impl.h
   types_ushort4.h
-  types_vector3.h
-  types_vector3_impl.h
   unique_ptr.h
   vector.h
   version.h
diff --git a/intern/cycles/util/types.h b/intern/cycles/util/types.h
index 031c2f7c4c1..26031d9e0fd 100644
--- a/intern/cycles/util/types.h
+++ b/intern/cycles/util/types.h
@@ -12,6 +12,7 @@
 
 #if !defined(__KERNEL_GPU__)
 #  include <stdint.h>
+#  include <stdio.h>
 #endif
 
 #include "util/defines.h"
@@ -70,6 +71,12 @@ ccl_device_inline bool is_power_of_two(size_t x)
 
 CCL_NAMESPACE_END
 
+/* Most GPU APIs matching native vector types, so we only need to implement them for
+ * CPU and oneAPI. */
+#if defined(__KERNEL_GPU__) && !defined(__KERNEL_ONEAPI__)
+#  define __KERNEL_NATIVE_VECTOR_TYPES__
+#endif
+
 /* Vectorized types declaration. */
 #include "util/types_uchar2.h"
 #include "util/types_uchar3.h"
@@ -90,8 +97,6 @@ CCL_NAMESPACE_END
 #include "util/types_float4.h"
 #include "util/types_float8.h"
 
-#include "util/types_vector3.h"
-
 /* Vectorized types implementation. */
 #include "util/types_uchar2_impl.h"
 #include "util/types_uchar3_impl.h"
@@ -110,8 +115,6 @@ CCL_NAMESPACE_END
 #include "util/types_float4_impl.h"
 #include "util/types_float8_impl.h"
 
-#include "util/types_vector3_impl.h"
-
 /* SSE types. */
 #ifndef __KERNEL_GPU__
 #  include "util/sseb.h"
diff --git a/intern/cycles/util/types_float2.h b/intern/cycles/util/types_float2.h
index 07b9ec0986b..f37aa1b4ad2 100644
--- a/intern/cycles/util/types_float2.h
+++ b/intern/cycles/util/types_float2.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT2_H__
-#define __UTIL_TYPES_FLOAT2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,18 +9,18 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct float2 {
   float x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline float operator[](int i) const;
   __forceinline float &operator[](int i);
+#  endif
 };
 
 ccl_device_inline float2 make_float2(float x, float y);
 ccl_device_inline void print_float2(const char *label, const float2 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT2_H__ */
diff --git a/intern/cycles/util/types_float2_impl.h b/intern/cycles/util/types_float2_impl.h
index 45fc90c52bd..9d1820fe17d 100644
--- a/intern/cycles/util/types_float2_impl.h
+++ b/intern/cycles/util/types_float2_impl.h
@@ -1,20 +1,16 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT2_IMPL_H__
-#define __UTIL_TYPES_FLOAT2_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline float float2::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -28,6 +24,7 @@ __forceinline float &float2::operator[](int i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline float2 make_float2(float x, float y)
 {
@@ -39,8 +36,6 @@ ccl_device_inline void print_float2(const char *label, const float2 &a)
 {
   printf("%s: %.8f %.8f\n", label, (double)a.x, (double)a.y);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT2_IMPL_H__ */
diff --git a/intern/cycles/util/types_float3.h b/intern/cycles/util/types_float3.h
index c7900acaa69..4e43e928657 100644
--- a/intern/cycles/util/types_float3.h
+++ b/intern/cycles/util/types_float3.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT3_H__
-#define __UTIL_TYPES_FLOAT3_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,28 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct ccl_try_align(16) float3
 {
-#  ifdef __KERNEL_SSE__
+#  ifdef __KERNEL_GPU__
+  /* Compact structure for GPU. */
+  float x, y, z;
+#  else
+  /* SIMD aligned structure for CPU. */
+#    ifdef __KERNEL_SSE__
   union {
     __m128 m128;
     struct {
       float x, y, z, w;
     };
   };
+#    else
+  float x, y, z, w;
+#    endif
+#  endif
 
+#  ifdef __KERNEL_SSE__
+  /* Convenient constructors and operators for SIMD, otherwise default is enough. */
   __forceinline float3();
   __forceinline float3(const float3 &a);
   __forceinline explicit float3(const __m128 &a);
@@ -29,18 +39,18 @@ struct ccl_try_align(16) float3
   __forceinline operator __m128 &();
 
   __forceinline float3 &operator=(const float3 &a);
-#  else  /* __KERNEL_SSE__ */
-  float x, y, z, w;
-#  endif /* __KERNEL_SSE__ */
+#  endif
 
+#  ifndef __KERNEL_GPU__
   __forceinline float operator[](int i) const;
   __forceinline float &operator[](int i);
+#  endif
 };
 
 ccl_device_inline float3 make_float3(float f);
 ccl_device_inline float3 make_float3(float x, float y, float z);
 ccl_device_inline void print_float3(const char *label, const float3 &a);
-#endif /* !defined(__KERNEL_GPU__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 /* Smaller float3 for storage. For math operations this must be converted to float3, so that on the
  * CPU SIMD instructions can be used. */
@@ -78,5 +88,3 @@ struct packed_float3 {
 static_assert(sizeof(packed_float3) == 12, "packed_float3 expected to be exactly 12 bytes");
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT3_H__ */
diff --git a/intern/cycles/util/types_float3_impl.h b/intern/cycles/util/types_float3_impl.h
index 2e6e864c8ea..cbd3f76dae4 100644
--- a/intern/cycles/util/types_float3_impl.h
+++ b/intern/cycles/util/types_float3_impl.h
@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT3_IMPL_H__
-#define __UTIL_TYPES_FLOAT3_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline float3::float3()
 {
@@ -45,6 +40,7 @@ __forceinline float3 &float3::operator=(const float3 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline float float3::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -58,23 +54,32 @@ __forceinline float &float3::operator[](int i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline float3 make_float3(float f)
 {
-#  ifdef __KERNEL_SSE__
-  float3 a(_mm_set1_ps(f));
+#  ifdef __KERNEL_GPU__
+  float3 a = {f, f, f};
 #  else
+#    ifdef __KERNEL_SSE__
+  float3 a(_mm_set1_ps(f));
+#    else
   float3 a = {f, f, f, f};
+#    endif
 #  endif
   return a;
 }
 
 ccl_device_inline float3 make_float3(float x, float y, float z)
 {
-#  ifdef __KERNEL_SSE__
-  float3 a(_mm_set_ps(0.0f, z, y, x));
+#  ifdef __KERNEL_GPU__
+  float3 a = {x, y, z};
 #  else
+#    ifdef __KERNEL_SSE__
+  float3 a(_mm_set_ps(0.0f, z, y, x));
+#    else
   float3 a = {x, y, z, 0.0f};
+#    endif
 #  endif
   return a;
 }
@@ -83,8 +88,6 @@ ccl_device_inline void print_float3(const char *label, const float3 &a)
 {
   printf("%s: %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z);
 }
-#endif /* !defined(__KERNEL_GPU__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT3_IMPL_H__ */
diff --git a/intern/cycles/util/types_float4.h b/intern/cycles/util/types_float4.h
index 27453bf39e4..6b301c47362 100644
--- a/intern/cycles/util/types_float4.h
+++ b/intern/cycles/util/types_float4.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT4_H__
-#define __UTIL_TYPES_FLOAT4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct int4;
 
 struct ccl_try_align(16) float4
@@ -35,16 +34,16 @@ struct ccl_try_align(16) float4
   float x, y, z, w;
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
   __forceinline float operator[](int i) const;
   __forceinline float &operator[](int i);
+#  endif
 };
 
 ccl_device_inline float4 make_float4(float f);
 ccl_device_inline float4 make_float4(float x, float y, float z, float w);
 ccl_device_inline float4 make_float4(const int4 &i);
 ccl_device_inline void print_float4(const char *label, const float4 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT4_H__ */
diff --git a/intern/cycles/util/types_float4_impl.h b/intern/cycles/util/types_float4_impl.h
index d7858f744e3..77b4fbff788 100644
--- a/intern/cycles/util/types_float4_impl.h
+++ b/intern/cycles/util/types_float4_impl.h
@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_FLOAT4_IMPL_H__
-#define __UTIL_TYPES_FLOAT4_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline float4::float4()
 {
@@ -41,6 +36,7 @@ __forceinline float4 &float4::operator=(const float4 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline float float4::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -54,6 +50,7 @@ __forceinline float &float4::operator[](int i)
   util_assert(i < 4);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline float4 make_float4(float f)
 {
@@ -89,8 +86,6 @@ ccl_device_inline void print_float4(const char *label, const float4 &a)
 {
   printf("%s: %.8f %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z, (double)a.w);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT4_IMPL_H__ */
diff --git a/intern/cycles/util/types_float8.h b/intern/cycles/util/types_float8.h
index bb9798932ac..29fd632f08e 100644
--- a/intern/cycles/util/types_float8.h
+++ b/intern/cycles/util/types_float8.h
@@ -2,8 +2,7 @@
  * Original code Copyright 2017, Intel Corporation
  * Modifications Copyright 2018-2022 Blender Foundation. */
 
-#ifndef __UTIL_TYPES_FLOAT8_H__
-#define __UTIL_TYPES_FLOAT8_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -12,7 +11,7 @@
 CCL_NAMESPACE_BEGIN
 
 /* float8 is a reserved type in Metal that has not been implemented. For
- * that reason this is named float8_t. */
+ * that reason this is named float8_t and not using native vector types. */
 
 #ifdef __KERNEL_GPU__
 struct float8_t
@@ -52,5 +51,3 @@ ccl_device_inline float8_t
 make_float8_t(float a, float b, float c, float d, float e, float f, float g, float h);
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT8_H__ */
diff --git a/intern/cycles/util/types_float8_impl.h b/intern/cycles/util/types_float8_impl.h
index 2ab464a791b..e8576cdaf70 100644
--- a/intern/cycles/util/types_float8_impl.h
+++ b/intern/cycles/util/types_float8_impl.h
@@ -2,17 +2,12 @@
  * Original code Copyright 2017, Intel Corporation
  * Modifications Copyright 2018-2022 Blender Foundation. */
 
-#ifndef __UTIL_TYPES_FLOAT8_IMPL_H__
-#define __UTIL_TYPES_FLOAT8_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
 #ifdef __KERNEL_AVX2__
@@ -83,5 +78,3 @@ make_float8_t(float a, float b, float c, float d, float e, float f, float g, flo
 }
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_FLOAT8_IMPL_H__ */
diff --git a/intern/cycles/util/types_int2.h b/intern/cycles/util/types_int2.h
index bf69cddc653..604713dffcd 100644
--- a/intern/cycles/util/types_int2.h
+++ b/intern/cycles/util/types_int2.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT2_H__
-#define __UTIL_TYPES_INT2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct int2 {
   int x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline int operator[](int i) const;
   __forceinline int &operator[](int i);
+#  endif
 };
 
 ccl_device_inline int2 make_int2(int x, int y);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT2_H__ */
diff --git a/intern/cycles/util/types_int2_impl.h b/intern/cycles/util/types_int2_impl.h
index 7bdc77369ee..f48c6f46729 100644
--- a/intern/cycles/util/types_int2_impl.h
+++ b/intern/cycles/util/types_int2_impl.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT2_IMPL_H__
-#define __UTIL_TYPES_INT2_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 int int2::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,14 +24,13 @@ int &int2::operator[](int i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline int2 make_int2(int x, int y)
 {
   int2 a = {x, y};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT2_IMPL_H__ */
diff --git a/intern/cycles/util/types_int3.h b/intern/cycles/util/types_int3.h
index f88ff22ac35..3196b158ee9 100644
--- a/intern/cycles/util/types_int3.h
+++ b/intern/cycles/util/types_int3.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT3_H__
-#define __UTIL_TYPES_INT3_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,10 +9,15 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct ccl_try_align(16) int3
 {
-#  ifdef __KERNEL_SSE__
+#  ifdef __KERNEL_GPU__
+  /* Compact structure on the GPU. */
+  int x, y, z;
+#  else
+  /* SIMD aligned structure for CPU. */
+#    ifdef __KERNEL_SSE__
   union {
     __m128i m128;
     struct {
@@ -29,19 +33,20 @@ struct ccl_try_align(16) int3
   __forceinline operator __m128i &();
 
   __forceinline int3 &operator=(const int3 &a);
-#  else  /* __KERNEL_SSE__ */
+#    else  /* __KERNEL_SSE__ */
   int x, y, z, w;
-#  endif /* __KERNEL_SSE__ */
+#    endif /* __KERNEL_SSE__ */
+#  endif
 
+#  ifndef __KERNEL_GPU__
   __forceinline int operator[](int i) const;
   __forceinline int &operator[](int i);
+#  endif
 };
 
 ccl_device_inline int3 make_int3(int i);
 ccl_device_inline int3 make_int3(int x, int y, int z);
 ccl_device_inline void print_int3(const char *label, const int3 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT3_H__ */
diff --git a/intern/cycles/util/types_int3_impl.h b/intern/cycles/util/types_int3_impl.h
index 1c49e97ad32..abc0f4e4309 100644
--- a/intern/cycles/util/types_int3_impl.h
+++ b/intern/cycles/util/types_int3_impl.h
@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT3_IMPL_H__
-#define __UTIL_TYPES_INT3_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline int3::int3()
 {
@@ -45,6 +40,7 @@ __forceinline int3 &int3::operator=(const int3 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline int int3::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -58,25 +54,33 @@ __forceinline int &int3::operator[](int i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline int3 make_int3(int i)
 {
-#  ifdef __KERNEL_SSE__
-  int3 a(_mm_set1_epi32(i));
+#  ifdef __KERNEL_GPU__
+  int3 a = {i, i, i};
 #  else
+#    ifdef __KERNEL_SSE__
+  int3 a(_mm_set1_epi32(i));
+#    else
   int3 a = {i, i, i, i};
+#    endif
 #  endif
   return a;
 }
 
 ccl_device_inline int3 make_int3(int x, int y, int z)
 {
-#  ifdef __KERNEL_SSE__
-  int3 a(_mm_set_epi32(0, z, y, x));
+#  ifdef __KERNEL_GPU__
+  int3 a = {x, y, z};
 #  else
+#    ifdef __KERNEL_SSE__
+  int3 a(_mm_set_epi32(0, z, y, x));
+#    else
   int3 a = {x, y, z, 0};
+#    endif
 #  endif
-
   return a;
 }
 
@@ -84,8 +88,6 @@ ccl_device_inline void print_int3(const char *label, const int3 &a)
 {
   printf("%s: %d %d %d\n", label, a.x, a.y, a.z);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT3_IMPL_H__ */
diff --git a/intern/cycles/util/types_int4.h b/intern/cycles/util/types_int4.h
index 9d557c01344..0ac53ffc09c 100644
--- a/intern/cycles/util/types_int4.h
+++ b/intern/cycles/util/types_int4.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT4_H__
-#define __UTIL_TYPES_INT4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 
 struct float3;
 struct float4;
@@ -37,8 +36,10 @@ struct ccl_try_align(16) int4
   int x, y, z, w;
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
   __forceinline int operator[](int i) const;
   __forceinline int &operator[](int i);
+#  endif
 };
 
 ccl_device_inline int4 make_int4(int i);
@@ -46,8 +47,6 @@ ccl_device_inline int4 make_int4(int x, int y, int z, int w);
 ccl_device_inline int4 make_int4(const float3 &f);
 ccl_device_inline int4 make_int4(const float4 &f);
 ccl_device_inline void print_int4(const char *label, const int4 &a);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT4_H__ */
diff --git a/intern/cycles/util/types_int4_impl.h b/intern/cycles/util/types_int4_impl.h
index 11e1ede6705..2aab29df941 100644
--- a/intern/cycles/util/types_int4_impl.h
+++ b/intern/cycles/util/types_int4_impl.h
@@ -1,20 +1,15 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_INT4_IMPL_H__
-#define __UTIL_TYPES_INT4_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
 #endif
 
-#ifndef __KERNEL_GPU__
-#  include <cstdio>
-#endif
-
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 #  ifdef __KERNEL_SSE__
 __forceinline int4::int4()
 {
@@ -45,6 +40,7 @@ __forceinline int4 &int4::operator=(const int4 &a)
 }
 #  endif /* __KERNEL_SSE__ */
 
+#  ifndef __KERNEL_GPU__
 __forceinline int int4::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -58,6 +54,7 @@ __forceinline int &int4::operator[](int i)
   util_assert(i < 4);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline int4 make_int4(int i)
 {
@@ -105,8 +102,6 @@ ccl_device_inline void print_int4(const char *label, const int4 &a)
 {
   printf("%s: %d %d %d %d\n", label, a.x, a.y, a.z, a.w);
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_INT4_IMPL_H__ */
diff --git a/intern/cycles/util/types_uchar2.h b/intern/cycles/util/types_uchar2.h
index 0b3c9bd0331..ce617248e6e 100644
--- a/intern/cycles/util/types_uchar2.h
+++ b/intern/cycles/util/types_uchar2.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UCHAR2_H__
-#define __UTIL_TYPES_UCHAR2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uchar2 {
   uchar x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uchar operator[](int i) const;
   __forceinline uchar &operator[](int i);
+#  endif
 };
 
 ccl_device_inline uchar2 make_uchar2(uchar x, uchar y);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UCHAR2_H__ */
diff --git a/intern/cycles/util/types_uchar2_impl.h b/intern/cycles/util/types_uchar2_impl.h
index a7254d5eaf2..9f3f3a4efb9 100644
--- a/intern/cycles/util/types_uchar2_impl.h
+++ b/intern/cycles/util/types_uchar2_impl.h
@@ -10,7 +10,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 uchar uchar2::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,13 +25,14 @@ uchar &uchar2::operator[](int i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uchar2 make_uchar2(uchar x, uchar y)
 {
   uchar2 a = {x, y};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/util/types_uchar3.h b/intern/cycles/util/types_uchar3.h
index fc213502ada..aed04c4775e 100644
--- a/intern/cycles/util/types_uchar3.h
+++ b/intern/cycles/util/types_uchar3.h
@@ -10,16 +10,18 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uchar3 {
   uchar x, y, z;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uchar operator[](int i) const;
   __forceinline uchar &operator[](int i);
+#  endif
 };
 
 ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/util/types_uchar3_impl.h b/intern/cycles/util/types_uchar3_impl.h
index 0c24ffb488a..83eb3c99b3c 100644
--- a/intern/cycles/util/types_uchar3_impl.h
+++ b/intern/cycles/util/types_uchar3_impl.h
@@ -10,7 +10,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 uchar uchar3::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,13 +25,14 @@ uchar &uchar3::operator[](int i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z)
 {
   uchar3 a = {x, y, z};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/util/types_uchar4.h b/intern/cycles/util/types_uchar4.h
index a2a2c945aaa..fb13a98875e 100644
--- a/intern/cycles/util/types_uchar4.h
+++ b/intern/cycles/util/types_uchar4.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UCHAR4_H__
-#define __UTIL_TYPES_UCHAR4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uchar4 {
   uchar x, y, z, w;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uchar operator[](int i) const;
   __forceinline uchar &operator[](int i);
+#  endif
 };
 
 ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UCHAR4_H__ */
diff --git a/intern/cycles/util/types_uchar4_impl.h b/intern/cycles/util/types_uchar4_impl.h
index 8ec6213a37d..244bb98f883 100644
--- a/intern/cycles/util/types_uchar4_impl.h
+++ b/intern/cycles/util/types_uchar4_impl.h
@@ -10,7 +10,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 uchar uchar4::operator[](int i) const
 {
   util_assert(i >= 0);
@@ -24,13 +25,14 @@ uchar &uchar4::operator[](int i)
   util_assert(i < 4);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w)
 {
   uchar4 a = {x, y, z, w};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/util/types_uint2.h b/intern/cycles/util/types_uint2.h
index faa0955f903..4d76b628088 100644
--- a/intern/cycles/util/types_uint2.h
+++ b/intern/cycles/util/types_uint2.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT2_H__
-#define __UTIL_TYPES_UINT2_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uint2 {
   uint x, y;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uint operator[](uint i) const;
   __forceinline uint &operator[](uint i);
+#  endif
 };
 
 ccl_device_inline uint2 make_uint2(uint x, uint y);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT2_H__ */
diff --git a/intern/cycles/util/types_uint2_impl.h b/intern/cycles/util/types_uint2_impl.h
index cac0ba6b531..b508aaf2543 100644
--- a/intern/cycles/util/types_uint2_impl.h
+++ b/intern/cycles/util/types_uint2_impl.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT2_IMPL_H__
-#define __UTIL_TYPES_UINT2_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline uint uint2::operator[](uint i) const
 {
   util_assert(i < 2);
@@ -22,14 +22,13 @@ __forceinline uint &uint2::operator[](uint i)
   util_assert(i < 2);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uint2 make_uint2(uint x, uint y)
 {
   uint2 a = {x, y};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT2_IMPL_H__ */
diff --git a/intern/cycles/util/types_uint3.h b/intern/cycles/util/types_uint3.h
index 3ff87bfc791..b1571716fc7 100644
--- a/intern/cycles/util/types_uint3.h
+++ b/intern/cycles/util/types_uint3.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT3_H__
-#define __UTIL_TYPES_UINT3_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uint3 {
   uint x, y, z;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uint operator[](uint i) const;
   __forceinline uint &operator[](uint i);
+#  endif
 };
 
 ccl_device_inline uint3 make_uint3(uint x, uint y, uint z);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT3_H__ */
diff --git a/intern/cycles/util/types_uint3_impl.h b/intern/cycles/util/types_uint3_impl.h
index 221883a1adb..d36c9f52de9 100644
--- a/intern/cycles/util/types_uint3_impl.h
+++ b/intern/cycles/util/types_uint3_impl.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT3_IMPL_H__
-#define __UTIL_TYPES_UINT3_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline uint uint3::operator[](uint i) const
 {
   util_assert(i < 3);
@@ -22,14 +22,13 @@ __forceinline uint &uint3::operator[](uint i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uint3 make_uint3(uint x, uint y, uint z)
 {
   uint3 a = {x, y, z};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT3_IMPL_H__ */
diff --git a/intern/cycles/util/types_uint4.h b/intern/cycles/util/types_uint4.h
index 504095b2383..4982b30f577 100644
--- a/intern/cycles/util/types_uint4.h
+++ b/intern/cycles/util/types_uint4.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT4_H__
-#define __UTIL_TYPES_UINT4_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,17 +9,17 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 struct uint4 {
   uint x, y, z, w;
 
+#  ifndef __KERNEL_GPU__
   __forceinline uint operator[](uint i) const;
   __forceinline uint &operator[](uint i);
+#  endif
 };
 
 ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w);
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT4_H__ */
diff --git a/intern/cycles/util/types_uint4_impl.h b/intern/cycles/util/types_uint4_impl.h
index d78db944a1f..1cfdb9e0992 100644
--- a/intern/cycles/util/types_uint4_impl.h
+++ b/intern/cycles/util/types_uint4_impl.h
@@ -1,8 +1,7 @@
 /* SPDX-License-Identifier: Apache-2.0
  * Copyright 2011-2022 Blender Foundation */
 
-#ifndef __UTIL_TYPES_UINT4_IMPL_H__
-#define __UTIL_TYPES_UINT4_IMPL_H__
+#pragma once
 
 #ifndef __UTIL_TYPES_H__
 #  error "Do not include this file directly, include util/types.h instead."
@@ -10,7 +9,8 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
+#  ifndef __KERNEL_GPU__
 __forceinline uint uint4::operator[](uint i) const
 {
   util_assert(i < 3);
@@ -22,14 +22,13 @@ __forceinline uint &uint4::operator[](uint i)
   util_assert(i < 3);
   return *(&x + i);
 }
+#  endif
 
 ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w)
 {
   uint4 a = {x, y, z, w};
   return a;
 }
-#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */
+#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
 
 CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_UINT4_IMPL_H__ */
diff --git a/intern/cycles/util/types_ushort4.h b/intern/cycles/util/types_ushort4.h
index 9a6e12095ba..aef36f63285 100644
--- a/intern/cycles/util/types_ushort4.h
+++ b/intern/cycles/util/types_ushort4.h
@@ -10,7 +10,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__)
+#ifndef __KERNEL_NATIVE_VECTOR_TYPES__
 
 struct ushort4 {
   uint16_t x, y, z, w;
diff --git a/intern/cycles/util/types_vector3.h b/intern/cycles/util/types_vector3.h
deleted file mode 100644
index 2e0d68e1bd0..00000000000
--- a/intern/cycles/util/types_vector3.h
+++ /dev/null
@@ -1,26 +0,0 @@
-/* SPDX-License-Identifier: Apache-2.0
- * Copyright 2011-2022 Blender Foundation */
-
-#ifndef __UTIL_TYPES_VECTOR3_H__
-#define __UTIL_TYPES_VECTOR3_H__
-
-#ifndef __UTIL_TYPES_H__
-#  error "Do not include this file directly, include util/types.h instead."
-#endif
-
-CCL_NAMESPACE_BEGIN
-
-#ifndef __KERNEL_GPU__
-template<typename T> class vector3 {
- public:
-  T x, y, z;
-
-  __forceinline vector3();
-  __forceinline vector3(const T &a);
-  __forceinline vector3(const T &x, const T &y, const T &z);
-};
-#endif /* __KERNEL_GPU__ */
-
-CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_VECTOR3_H__ */
diff --git a/intern/cycles/util/types_vector3_impl.h b/intern/cycles/util/types_vector3_impl.h
deleted file mode 100644
index a765780e2d3..00000000000
--- a/intern/cycles/util/types_vector3_impl.h
+++ /dev/null
@@ -1,30 +0,0 @@
-/* SPDX-License-Identifier: Apache-2.0
- * Copyright 2011-2022 Blender Foundation */
-
-#ifndef __UTIL_TYPES_VECTOR3_IMPL_H__
-#define __UTIL_TYPES_VECTOR3_IMPL_H__
-
-#ifndef __UTIL_TYPES_H__
-#  error "Do not include this file directly, include util/types.h instead."
-#endif
-
-CCL_NAMESPACE_BEGIN
-
-#ifndef __KERNEL_GPU__
-template<typename T> ccl_always_inline vector3<T>::vector3()
-{
-}
-
-template<typename T> ccl_always_inline vector3<T>::vector3(const T &a) : x(a), y(a), z(a)
-{
-}
-
-template<typename T>
-ccl_always_inline vector3<T>::vector3(const T &x, const T &y, const T &z) : x(x), y(y), z(z)
-{
-}
-#endif /* __KERNEL_GPU__ */
-
-CCL_NAMESPACE_END
-
-#endif /* __UTIL_TYPES_VECTOR3_IMPL_H__ */
-- 
cgit v1.2.3