Cleanup: renaming and consistency for kernel data

* Rename "texture" to "data array". This has not used textures for a long time,
  there are just global memory arrays now. (On old CUDA GPUs there was a cache
  for textures but not global memory, so we used to put all data in textures.)
* For CUDA and HIP, put globals in KernelParams struct like other devices.
* Drop __ prefix for data array names, no possibility for naming conflict now that
  these are in a struct.

1 files changed, 16 insertions, 10 deletions

diff --git a/intern/cycles/kernel/device/cuda/globals.h b/intern/cycles/kernel/device/cuda/globals.h
index e77fcd2b424..f5f7bcf58ee 100644
--- a/intern/cycles/kernel/device/cuda/globals.h
+++ b/intern/cycles/kernel/device/cuda/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) const __constant__ __device__ type *name;
-#include "kernel/textures.h"
+struct KernelParamsCUDA {
+  /* 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__ KernelParamsCUDA 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