1 files changed, 16 insertions, 10 deletions
diff --git a/intern/cycles/kernel/device/hip/globals.h b/intern/cycles/kernel/device/hip/globals.h
index 50f117038a2..3a334b21a9e 100644
--- a/intern/cycles/kernel/device/hip/globals.h
+++ b/intern/cycles/kernel/device/hip/globals.h
@@ -20,18 +20,24 @@ struct KernelGlobalsGPU {
 };
 typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
 
-/* Global scene data and textures */
-__constant__ KernelData __data;
-#define KERNEL_TEX(type, name) __attribute__((used)) const __constant__ __device__ type *name;
-#include "kernel/textures.h"
+struct KernelParamsHIP {
+  /* Global scene data and textures */
+  KernelData data;
+#define KERNEL_DATA_ARRAY(type, name) const type *name;
+#include "kernel/data_arrays.h"
+
+  /* Integrator state */
+  IntegratorStateGPU integrator_state;
+};
 
-/* Integrator state */
-__constant__ IntegratorStateGPU __integrator_state;
+#ifdef __KERNEL_GPU__
+__constant__ KernelParamsHIP kernel_params;
+#endif
 
 /* Abstraction macros */
-#define kernel_data __data
-#define kernel_tex_fetch(t, index) t[(index)]
-#define kernel_tex_array(t) (t)
-#define kernel_integrator_state __integrator_state
+#define kernel_data kernel_params.data
+#define kernel_data_fetch(name, index) kernel_params.name[(index)]
+#define kernel_data_array(name) (kernel_params.name)
+#define kernel_integrator_state kernel_params.integrator_state
 
 CCL_NAMESPACE_END