ilusha/blender
1
0
Fork 0
mirror of https://github.com/blender/blender synced 2026-09-29 04:37:17 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Refactor: Cycles: Add writeable global GPU data arrays

Browse source 
This will be used for tile requests in the texture cache.

Pull Request: https://projects.blender.org/blender/blender/pulls/154913
This commit is contained in:
Brecht Van Lommel 2026-02-14 23:58:52 +01:00
parent af30c5b076
commit bac976d38f
8 changed files with 26 additions and 5 deletions
Download patch file Download diff file Expand all files Collapse all files

4
intern/cycles/kernel/data_arrays.h
View file

@ -7,6 +7,9 @@
#ifndef KERNEL_DATA_ARRAY
# define KERNEL_DATA_ARRAY(type, name)
#endif
#ifndef KERNEL_DATA_ARRAY_WRITABLE
# define KERNEL_DATA_ARRAY_WRITABLE(type, name) KERNEL_DATA_ARRAY(type, name)
#endif
/* BVH2, not used for OptiX or Embree. */
KERNEL_DATA_ARRAY(float4, bvh_nodes)
@ -92,3 +95,4 @@ KERNEL_DATA_ARRAY(KernelImageUDIM, image_texture_udims)
KERNEL_DATA_ARRAY(KernelImageInfo, image_info)
#undef KERNEL_DATA_ARRAY
#undef KERNEL_DATA_ARRAY_WRITABLE

9
intern/cycles/kernel/device/cpu/globals.h
View file

@ -34,13 +34,19 @@ template<typename T> struct kernel_array {
return data[index];
}
ccl_always_inline void write(const int index, const T &value) const
{
data[index] = value;
}
T *data = nullptr;
int width = 0;
};
/* Constant globals shared between all threads. */
struct KernelGlobalsCPU {
#define KERNEL_DATA_ARRAY(type, name) kernel_array<type> name;
#define KERNEL_DATA_ARRAY(type, name) kernel_array<const type> name;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) kernel_array<type> name;
#include "kernel/data_arrays.h"
KernelData data = {};
@ -93,6 +99,7 @@ using KernelGlobals = const ThreadKernelGlobalsCPU *;
/* Abstraction macros */
#define kernel_data_fetch(name, index) (kg->name.fetch(index))
#define kernel_data_write(name, index, value) (kg->name.write(index, value))
#define kernel_data_array(name) (kg->name.data)
#define kernel_data (kg->data)
#if defined(WITH_PATH_GUIDING)

2
intern/cycles/kernel/device/cuda/globals.h
View file

@ -27,6 +27,7 @@ struct KernelParamsCUDA {
/* Global scene data and textures */
KernelData data;
#define KERNEL_DATA_ARRAY(type, name) const type *name;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) type *name;
#include "kernel/data_arrays.h"
/* Integrator state */
@ -40,6 +41,7 @@ __constant__ KernelParamsCUDA kernel_params;
/* Abstraction macros */
#define kernel_data kernel_params.data
#define kernel_data_fetch(name, index) kernel_params.name[(index)]
#define kernel_data_write(name, index, value) kernel_params.name[(index)] = (value)
#define kernel_data_array(name) (kernel_params.name)
#define kernel_integrator_state kernel_params.integrator_state

2
intern/cycles/kernel/device/hip/globals.h
View file

@ -27,6 +27,7 @@ struct KernelParamsHIP {
/* Global scene data and textures */
KernelData data;
#define KERNEL_DATA_ARRAY(type, name) const type *name;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) type *name;
#include "kernel/data_arrays.h"
/* Integrator state */
@ -40,6 +41,7 @@ __constant__ KernelParamsHIP kernel_params;
/* Abstraction macros */
#define kernel_data kernel_params.data
#define kernel_data_fetch(name, index) kernel_params.name[(index)]
#define kernel_data_write(name, index, value) kernel_params.name[(index)] = (value)
#define kernel_data_array(name) (kernel_params.name)
#define kernel_integrator_state kernel_params.integrator_state

2
intern/cycles/kernel/device/hiprt/globals.h
View file

@ -52,6 +52,7 @@ using KernelGlobals = ccl_global KernelGlobalsGPU *ccl_restrict;
struct KernelParamsHIPRT {
KernelData data;
#define KERNEL_DATA_ARRAY(type, name) const type *name;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) type *name;
KERNEL_DATA_ARRAY(int, user_instance_id)
KERNEL_DATA_ARRAY(uint64_t, blas_ptr)
KERNEL_DATA_ARRAY(int2, custom_prim_info)
@ -145,6 +146,7 @@ typedef hiprtEmptyInstanceStack Instance_Stack;
/* Abstraction macros */
#define kernel_data kernel_params.data
#define kernel_data_fetch(name, index) kernel_params.name[(index)]
#define kernel_data_write(name, index, value) kernel_params.name[(index)] = (value)
#define kernel_data_array(name) (kernel_params.name)
#define kernel_integrator_state kernel_params.integrator_state

5
intern/cycles/kernel/device/metal/globals.h
View file

@ -16,9 +16,9 @@ CCL_NAMESPACE_BEGIN
struct KernelParamsMetal {
#define KERNEL_DATA_ARRAY(type, name) ccl_global const type *name;
#define KERNEL_DATA_ARRAY(type, name) const ccl_global type *name;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) ccl_global type *name;
#include "kernel/data_arrays.h"
#undef KERNEL_DATA_ARRAY
const IntegratorStateGPU integrator_state;
const KernelData data;
@ -33,6 +33,7 @@ using KernelGlobals = const ccl_global KernelGlobalsGPU *ccl_restrict;
/* Abstraction macros */
#define kernel_data launch_params_metal.data
#define kernel_data_fetch(name, index) launch_params_metal.name[index]
#define kernel_data_write(name, index, value) launch_params_metal.name[index] = (value)
#define kernel_data_array(name) launch_params_metal.name
#define kernel_integrator_state launch_params_metal.integrator_state

5
intern/cycles/kernel/device/oneapi/globals.h
View file

@ -26,8 +26,8 @@ struct IntegratorQueueCounter;
struct KernelGlobalsGPU {
#define KERNEL_DATA_ARRAY(type, name) const type *__##name = nullptr;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) type *__##name = nullptr;
#include "kernel/data_arrays.h"
#undef KERNEL_DATA_ARRAY
IntegratorStateGPU *integrator_state;
const KernelData *__data;
sycl::kernel_handler kernel_handler;
@ -40,7 +40,8 @@ using KernelGlobals = ccl_global KernelGlobalsGPU *ccl_restrict;
/* data lookup defines */
#define kernel_data_fetch(name, index) __##name[index]
#define kernel_data_fetch(name, index) __##name[(index)]
#define kernel_data_write(name, index, value) __##name[(index)] = (value)
#define kernel_data_array(name) __##name
CCL_NAMESPACE_END

2
intern/cycles/kernel/device/optix/globals.h
View file

@ -38,6 +38,7 @@ struct KernelParamsOptiX {
/* Global scene data and textures */
KernelData data;
#define KERNEL_DATA_ARRAY(type, name) const type *name;
#define KERNEL_DATA_ARRAY_WRITABLE(type, name) type *name;
#include "kernel/data_arrays.h"
/* Integrator state */
@ -58,6 +59,7 @@ extern "C"
#define kernel_data kernel_params.data
#define kernel_data_array(name) kernel_params.name
#define kernel_data_fetch(name, index) kernel_params.name[(index)]
#define kernel_data_write(name, index, value) kernel_params.name[(index)] = (value)
#define kernel_integrator_state kernel_params.integrator_state
CCL_NAMESPACE_END