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
Referenced by commit 4b14b33ea8, Cycles: use packed float3 back for oneAPI
2022-07-28 19:57:30 +02:00 · 2022-07-28 19:57:30 +02:00 · 79ab76e156 · 2023-02-13 23:39:48 +01:00
parent d094a3722c
commit 79ab76e156
34 changed files with 175 additions and 290 deletions
--- 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
--- 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
--- 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))
--- 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
--- 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"
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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

--- 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

--- 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

--- 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__ */
--- 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

--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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__ */
--- 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;
--- a/intern/cycles/util/types_vector3.h
+++ b/intern/cycles/util/types_vector3.h
@ -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__ */
--- a/intern/cycles/util/types_vector3_impl.h
+++ b/intern/cycles/util/types_vector3_impl.h
@ -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__ */