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.
This commit is contained in:
@@ -250,7 +250,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
|
||||
*isect = current_isect;
|
||||
/* Only primitives from volume object. */
|
||||
uint tri_object = isect->object;
|
||||
int object_flag = kernel_tex_fetch(__object_flag, tri_object);
|
||||
int object_flag = kernel_data_fetch(object_flag, tri_object);
|
||||
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
--ctx->num_hits;
|
||||
}
|
||||
|
@@ -51,7 +51,7 @@
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
CPUDevice::CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_)
|
||||
: Device(info_, stats_, profiler_), texture_info(this, "__texture_info", MEM_GLOBAL)
|
||||
: Device(info_, stats_, profiler_), texture_info(this, "texture_info", MEM_GLOBAL)
|
||||
{
|
||||
/* Pick any kernel, all of them are supposed to have same level of microarchitecture
|
||||
* optimization. */
|
||||
@@ -192,7 +192,7 @@ device_ptr CPUDevice::mem_alloc_sub_ptr(device_memory &mem, size_t offset, size_
|
||||
void CPUDevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
{
|
||||
#ifdef WITH_EMBREE
|
||||
if (strcmp(name, "__data") == 0) {
|
||||
if (strcmp(name, "data") == 0) {
|
||||
assert(size <= sizeof(KernelData));
|
||||
|
||||
// Update scene handle (since it is different for each device on multi devices)
|
||||
|
@@ -23,6 +23,8 @@
|
||||
# include "util/types.h"
|
||||
# include "util/windows.h"
|
||||
|
||||
# include "kernel/device/cuda/globals.h"
|
||||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
class CUDADevice;
|
||||
@@ -51,7 +53,7 @@ void CUDADevice::set_error(const string &error)
|
||||
}
|
||||
|
||||
CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||
: Device(info, stats, profiler), texture_info(this, "__texture_info", MEM_GLOBAL)
|
||||
: Device(info, stats, profiler), texture_info(this, "texture_info", MEM_GLOBAL)
|
||||
{
|
||||
first_error = true;
|
||||
|
||||
@@ -900,9 +902,19 @@ void CUDADevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
CUdeviceptr mem;
|
||||
size_t bytes;
|
||||
|
||||
cuda_assert(cuModuleGetGlobal(&mem, &bytes, cuModule, name));
|
||||
// assert(bytes == size);
|
||||
cuda_assert(cuMemcpyHtoD(mem, host, size));
|
||||
cuda_assert(cuModuleGetGlobal(&mem, &bytes, cuModule, "kernel_params"));
|
||||
assert(bytes == sizeof(KernelParamsCUDA));
|
||||
|
||||
/* Update data storage pointers in launch parameters. */
|
||||
# define KERNEL_DATA_ARRAY(data_type, data_name) \
|
||||
if (strcmp(name, #data_name) == 0) { \
|
||||
cuda_assert(cuMemcpyHtoD(mem + offsetof(KernelParamsCUDA, data_name), host, size)); \
|
||||
return; \
|
||||
}
|
||||
KERNEL_DATA_ARRAY(KernelData, data)
|
||||
KERNEL_DATA_ARRAY(IntegratorStateGPU, integrator_state)
|
||||
# include "kernel/data_arrays.h"
|
||||
# undef KERNEL_DATA_ARRAY
|
||||
}
|
||||
|
||||
void CUDADevice::global_alloc(device_memory &mem)
|
||||
@@ -926,7 +938,6 @@ void CUDADevice::tex_alloc(device_texture &mem)
|
||||
{
|
||||
CUDAContextScope scope(this);
|
||||
|
||||
string bind_name = mem.name;
|
||||
size_t dsize = datatype_size(mem.data_type);
|
||||
size_t size = mem.memory_size();
|
||||
|
||||
|
@@ -24,6 +24,8 @@
|
||||
# include "util/types.h"
|
||||
# include "util/windows.h"
|
||||
|
||||
# include "kernel/device/hip/globals.h"
|
||||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
class HIPDevice;
|
||||
@@ -52,7 +54,7 @@ void HIPDevice::set_error(const string &error)
|
||||
}
|
||||
|
||||
HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||
: Device(info, stats, profiler), texture_info(this, "__texture_info", MEM_GLOBAL)
|
||||
: Device(info, stats, profiler), texture_info(this, "texture_info", MEM_GLOBAL)
|
||||
{
|
||||
first_error = true;
|
||||
|
||||
@@ -856,8 +858,19 @@ void HIPDevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
hipDeviceptr_t mem;
|
||||
size_t bytes;
|
||||
|
||||
hip_assert(hipModuleGetGlobal(&mem, &bytes, hipModule, name));
|
||||
hip_assert(hipMemcpyHtoD(mem, host, size));
|
||||
hip_assert(hipModuleGetGlobal(&mem, &bytes, hipModule, "kernel_params"));
|
||||
assert(bytes == sizeof(KernelParamsHIP));
|
||||
|
||||
/* Update data storage pointers in launch parameters. */
|
||||
# define KERNEL_DATA_ARRAY(data_type, data_name) \
|
||||
if (strcmp(name, #data_name) == 0) { \
|
||||
hip_assert(hipMemcpyHtoD(mem + offsetof(KernelParamsHIP, data_name), host, size)); \
|
||||
return; \
|
||||
}
|
||||
KERNEL_DATA_ARRAY(KernelData, data)
|
||||
KERNEL_DATA_ARRAY(IntegratorStateGPU, integrator_state)
|
||||
# include "kernel/data_arrays.h"
|
||||
# undef KERNEL_DATA_ARRAY
|
||||
}
|
||||
|
||||
void HIPDevice::global_alloc(device_memory &mem)
|
||||
@@ -881,7 +894,6 @@ void HIPDevice::tex_alloc(device_texture &mem)
|
||||
{
|
||||
HIPContextScope scope(this);
|
||||
|
||||
string bind_name = mem.name;
|
||||
size_t dsize = datatype_size(mem.data_type);
|
||||
size_t size = mem.memory_size();
|
||||
|
||||
|
@@ -350,7 +350,7 @@ template<typename T> class device_only_memory : public device_memory {
|
||||
*
|
||||
* When using memory type MEM_GLOBAL, a pointer to this memory will be
|
||||
* automatically attached to kernel globals, using the provided name
|
||||
* matching an entry in kernel_textures.h. */
|
||||
* matching an entry in kernel/data_arrays.h. */
|
||||
|
||||
template<typename T> class device_vector : public device_memory {
|
||||
public:
|
||||
|
@@ -35,7 +35,7 @@ void MetalDevice::set_error(const string &error)
|
||||
}
|
||||
|
||||
MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||
: Device(info, stats, profiler), texture_info(this, "__texture_info", MEM_GLOBAL)
|
||||
: Device(info, stats, profiler), texture_info(this, "texture_info", MEM_GLOBAL)
|
||||
{
|
||||
mtlDevId = info.num;
|
||||
|
||||
@@ -625,7 +625,7 @@ device_ptr MetalDevice::mem_alloc_sub_ptr(device_memory &mem, size_t offset, siz
|
||||
|
||||
void MetalDevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
{
|
||||
if (strcmp(name, "__data") == 0) {
|
||||
if (strcmp(name, "data") == 0) {
|
||||
assert(size == sizeof(KernelData));
|
||||
memcpy((uint8_t *)&launch_params + offsetof(KernelParamsMetal, data), host, size);
|
||||
return;
|
||||
@@ -646,19 +646,19 @@ void MetalDevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
};
|
||||
|
||||
/* Update data storage pointers in launch parameters. */
|
||||
if (strcmp(name, "__integrator_state") == 0) {
|
||||
if (strcmp(name, "integrator_state") == 0) {
|
||||
/* IntegratorStateGPU is contiguous pointers */
|
||||
const size_t pointer_block_size = sizeof(IntegratorStateGPU);
|
||||
update_launch_pointers(
|
||||
offsetof(KernelParamsMetal, __integrator_state), host, size, pointer_block_size);
|
||||
offsetof(KernelParamsMetal, integrator_state), host, size, pointer_block_size);
|
||||
}
|
||||
# define KERNEL_TEX(data_type, tex_name) \
|
||||
# define KERNEL_DATA_ARRAY(data_type, tex_name) \
|
||||
else if (strcmp(name, #tex_name) == 0) \
|
||||
{ \
|
||||
update_launch_pointers(offsetof(KernelParamsMetal, tex_name), host, size, size); \
|
||||
}
|
||||
# include "kernel/textures.h"
|
||||
# undef KERNEL_TEX
|
||||
# include "kernel/data_arrays.h"
|
||||
# undef KERNEL_DATA_ARRAY
|
||||
}
|
||||
|
||||
void MetalDevice::global_alloc(device_memory &mem)
|
||||
|
@@ -358,7 +358,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
|
||||
|
||||
/* Prepare any non-pointer (i.e. plain-old-data) KernelParamsMetal data */
|
||||
/* The plain-old-data is contiguous, continuing to the end of KernelParamsMetal */
|
||||
size_t plain_old_launch_data_offset = offsetof(KernelParamsMetal, __integrator_state) +
|
||||
size_t plain_old_launch_data_offset = offsetof(KernelParamsMetal, integrator_state) +
|
||||
sizeof(IntegratorStateGPU);
|
||||
size_t plain_old_launch_data_size = sizeof(KernelParamsMetal) - plain_old_launch_data_offset;
|
||||
memcpy(init_arg_buffer + globals_offsets + plain_old_launch_data_offset,
|
||||
@@ -415,7 +415,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
|
||||
}
|
||||
|
||||
/* this relies on IntegratorStateGPU layout being contiguous device_ptrs */
|
||||
const size_t pointer_block_end = offsetof(KernelParamsMetal, __integrator_state) +
|
||||
const size_t pointer_block_end = offsetof(KernelParamsMetal, integrator_state) +
|
||||
sizeof(IntegratorStateGPU);
|
||||
for (size_t offset = 0; offset < pointer_block_end; offset += sizeof(device_ptr)) {
|
||||
int pointer_index = int(offset / sizeof(device_ptr));
|
||||
|
@@ -246,7 +246,7 @@ OptiXDevice::Denoiser::Denoiser(OptiXDevice *device)
|
||||
OptiXDevice::OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||
: CUDADevice(info, stats, profiler),
|
||||
sbt_data(this, "__sbt", MEM_READ_ONLY),
|
||||
launch_params(this, "__params", false),
|
||||
launch_params(this, "kernel_params", false),
|
||||
denoiser_(this)
|
||||
{
|
||||
/* Make the CUDA context current. */
|
||||
@@ -421,7 +421,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
|
||||
pipeline_options.numPayloadValues = 8;
|
||||
pipeline_options.numAttributeValues = 2; /* u, v */
|
||||
pipeline_options.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
|
||||
pipeline_options.pipelineLaunchParamsVariableName = "__params"; /* See globals.h */
|
||||
pipeline_options.pipelineLaunchParamsVariableName = "kernel_params"; /* See globals.h */
|
||||
|
||||
pipeline_options.usesPrimitiveTypeFlags = OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE;
|
||||
if (kernel_features & KERNEL_FEATURE_HAIR) {
|
||||
@@ -2042,7 +2042,7 @@ void OptiXDevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
/* Set constant memory for CUDA module. */
|
||||
CUDADevice::const_copy_to(name, host, size);
|
||||
|
||||
if (strcmp(name, "__data") == 0) {
|
||||
if (strcmp(name, "data") == 0) {
|
||||
assert(size <= sizeof(KernelData));
|
||||
|
||||
/* Update traversable handle (since it is different for each device on multi devices). */
|
||||
@@ -2054,14 +2054,14 @@ void OptiXDevice::const_copy_to(const char *name, void *host, size_t size)
|
||||
}
|
||||
|
||||
/* Update data storage pointers in launch parameters. */
|
||||
# define KERNEL_TEX(data_type, tex_name) \
|
||||
if (strcmp(name, #tex_name) == 0) { \
|
||||
update_launch_params(offsetof(KernelParamsOptiX, tex_name), host, size); \
|
||||
# define KERNEL_DATA_ARRAY(data_type, data_name) \
|
||||
if (strcmp(name, #data_name) == 0) { \
|
||||
update_launch_params(offsetof(KernelParamsOptiX, data_name), host, size); \
|
||||
return; \
|
||||
}
|
||||
KERNEL_TEX(IntegratorStateGPU, __integrator_state)
|
||||
# include "kernel/textures.h"
|
||||
# undef KERNEL_TEX
|
||||
KERNEL_DATA_ARRAY(IntegratorStateGPU, integrator_state)
|
||||
# include "kernel/data_arrays.h"
|
||||
# undef KERNEL_DATA_ARRAY
|
||||
}
|
||||
|
||||
void OptiXDevice::update_launch_params(size_t offset, void *data, size_t data_size)
|
||||
|
@@ -239,7 +239,7 @@ void PathTraceWorkGPU::init_execution()
|
||||
|
||||
/* Copy to device side struct in constant memory. */
|
||||
device_->const_copy_to(
|
||||
"__integrator_state", &integrator_state_gpu_, sizeof(integrator_state_gpu_));
|
||||
"integrator_state", &integrator_state_gpu_, sizeof(integrator_state_gpu_));
|
||||
}
|
||||
|
||||
void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
|
||||
|
@@ -267,8 +267,8 @@ set(SRC_KERNEL_UTIL_HEADERS
|
||||
)
|
||||
|
||||
set(SRC_KERNEL_TYPES_HEADERS
|
||||
data_arrays.h
|
||||
tables.h
|
||||
textures.h
|
||||
types.h
|
||||
)
|
||||
|
||||
|
@@ -452,7 +452,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
|
||||
|
||||
# ifdef __EMBREE__
|
||||
if (kernel_data.bvh.scene) {
|
||||
const bool has_bvh = !(kernel_tex_fetch(__object_flag, local_object) &
|
||||
const bool has_bvh = !(kernel_data_fetch(object_flag, local_object) &
|
||||
SD_OBJECT_TRANSFORM_APPLIED);
|
||||
CCLIntersectContext ctx(
|
||||
kg, has_bvh ? CCLIntersectContext::RAY_SSS : CCLIntersectContext::RAY_LOCAL);
|
||||
|
@@ -146,14 +146,14 @@ ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,
|
||||
|
||||
const bool is_hair = hit->geomID & 1;
|
||||
if (is_hair) {
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, isect->prim);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, isect->prim);
|
||||
isect->type = segment.type;
|
||||
isect->prim = segment.prim;
|
||||
isect->u = hit->u;
|
||||
isect->v = hit->v;
|
||||
}
|
||||
else {
|
||||
isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
|
||||
isect->type = kernel_data_fetch(objects, isect->object).primitive_type;
|
||||
isect->u = 1.0f - hit->v - hit->u;
|
||||
isect->v = hit->u;
|
||||
}
|
||||
@@ -170,7 +170,7 @@ ccl_device_inline void kernel_embree_convert_sss_hit(
|
||||
isect->prim = hit->primID +
|
||||
(intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
|
||||
isect->object = object;
|
||||
isect->type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
isect->type = kernel_data_fetch(objects, object).primitive_type;
|
||||
}
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -41,7 +41,7 @@ ccl_device_inline
|
||||
|
||||
/* traversal variables in registers */
|
||||
int stack_ptr = 0;
|
||||
int node_addr = kernel_tex_fetch(__object_node, local_object);
|
||||
int node_addr = kernel_data_fetch(object_node, local_object);
|
||||
|
||||
/* ray parameters in registers */
|
||||
float3 P = ray->P;
|
||||
@@ -55,7 +55,7 @@ ccl_device_inline
|
||||
}
|
||||
kernel_assert((local_isect == NULL) == (max_hits == 0));
|
||||
|
||||
const int object_flag = kernel_tex_fetch(__object_flag, local_object);
|
||||
const int object_flag = kernel_data_fetch(object_flag, local_object);
|
||||
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
|
||||
#if BVH_FEATURE(BVH_MOTION)
|
||||
Transform ob_itfm;
|
||||
@@ -73,7 +73,7 @@ ccl_device_inline
|
||||
while (node_addr >= 0 && node_addr != ENTRYPOINT_SENTINEL) {
|
||||
int node_addr_child1, traverse_mask;
|
||||
float dist[2];
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
|
||||
traverse_mask = NODE_INTERSECT(kg,
|
||||
P,
|
||||
@@ -117,7 +117,7 @@ ccl_device_inline
|
||||
|
||||
/* if node is leaf, fetch triangle list */
|
||||
if (node_addr < 0) {
|
||||
float4 leaf = kernel_tex_fetch(__bvh_leaf_nodes, (-node_addr - 1));
|
||||
float4 leaf = kernel_data_fetch(bvh_leaf_nodes, (-node_addr - 1));
|
||||
int prim_addr = __float_as_int(leaf.x);
|
||||
|
||||
const int prim_addr2 = __float_as_int(leaf.y);
|
||||
@@ -132,18 +132,18 @@ ccl_device_inline
|
||||
case PRIMITIVE_TRIANGLE: {
|
||||
/* intersect ray against primitive */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
|
||||
/* Only intersect with matching object, for instanced objects we
|
||||
* already know we are only intersecting the right object. */
|
||||
if (object == OBJECT_NONE) {
|
||||
if (kernel_tex_fetch(__prim_object, prim_addr) != local_object) {
|
||||
if (kernel_data_fetch(prim_object, prim_addr) != local_object) {
|
||||
continue;
|
||||
}
|
||||
}
|
||||
|
||||
/* Skip self intersection. */
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self_local(ray->self, prim)) {
|
||||
continue;
|
||||
}
|
||||
@@ -167,18 +167,18 @@ ccl_device_inline
|
||||
case PRIMITIVE_MOTION_TRIANGLE: {
|
||||
/* intersect ray against primitive */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
|
||||
/* Only intersect with matching object, for instanced objects we
|
||||
* already know we are only intersecting the right object. */
|
||||
if (object == OBJECT_NONE) {
|
||||
if (kernel_tex_fetch(__prim_object, prim_addr) != local_object) {
|
||||
if (kernel_data_fetch(prim_object, prim_addr) != local_object) {
|
||||
continue;
|
||||
}
|
||||
}
|
||||
|
||||
/* Skip self intersection. */
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self_local(ray->self, prim)) {
|
||||
continue;
|
||||
}
|
||||
|
@@ -9,9 +9,9 @@ ccl_device_forceinline Transform bvh_unaligned_node_fetch_space(KernelGlobals kg
|
||||
{
|
||||
Transform space;
|
||||
const int child_addr = node_addr + child * 3;
|
||||
space.x = kernel_tex_fetch(__bvh_nodes, child_addr + 1);
|
||||
space.y = kernel_tex_fetch(__bvh_nodes, child_addr + 2);
|
||||
space.z = kernel_tex_fetch(__bvh_nodes, child_addr + 3);
|
||||
space.x = kernel_data_fetch(bvh_nodes, child_addr + 1);
|
||||
space.y = kernel_data_fetch(bvh_nodes, child_addr + 2);
|
||||
space.z = kernel_data_fetch(bvh_nodes, child_addr + 3);
|
||||
return space;
|
||||
}
|
||||
|
||||
@@ -26,11 +26,11 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
|
||||
|
||||
/* fetch node data */
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
#endif
|
||||
float4 node0 = kernel_tex_fetch(__bvh_nodes, node_addr + 1);
|
||||
float4 node1 = kernel_tex_fetch(__bvh_nodes, node_addr + 2);
|
||||
float4 node2 = kernel_tex_fetch(__bvh_nodes, node_addr + 3);
|
||||
float4 node0 = kernel_data_fetch(bvh_nodes, node_addr + 1);
|
||||
float4 node1 = kernel_data_fetch(bvh_nodes, node_addr + 2);
|
||||
float4 node2 = kernel_data_fetch(bvh_nodes, node_addr + 3);
|
||||
|
||||
/* intersect ray against child nodes */
|
||||
float c0lox = (node0.x - P.x) * idir.x;
|
||||
@@ -100,7 +100,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
|
||||
{
|
||||
int mask = 0;
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
#endif
|
||||
if (bvh_unaligned_node_intersect_child(kg, P, dir, t, node_addr, 0, &dist[0])) {
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
@@ -130,7 +130,7 @@ ccl_device_forceinline int bvh_node_intersect(KernelGlobals kg,
|
||||
const uint visibility,
|
||||
float dist[2])
|
||||
{
|
||||
float4 node = kernel_tex_fetch(__bvh_nodes, node_addr);
|
||||
float4 node = kernel_data_fetch(bvh_nodes, node_addr);
|
||||
if (__float_as_uint(node.x) & PATH_RAY_NODE_UNALIGNED) {
|
||||
return bvh_unaligned_node_intersect(kg, P, dir, idir, t, node_addr, visibility, dist);
|
||||
}
|
||||
|
@@ -80,7 +80,7 @@ ccl_device_inline
|
||||
while (node_addr >= 0 && node_addr != ENTRYPOINT_SENTINEL) {
|
||||
int node_addr_child1, traverse_mask;
|
||||
float dist[2];
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
|
||||
traverse_mask = NODE_INTERSECT(kg,
|
||||
P,
|
||||
@@ -124,7 +124,7 @@ ccl_device_inline
|
||||
|
||||
/* if node is leaf, fetch triangle list */
|
||||
if (node_addr < 0) {
|
||||
float4 leaf = kernel_tex_fetch(__bvh_leaf_nodes, (-node_addr - 1));
|
||||
float4 leaf = kernel_data_fetch(bvh_leaf_nodes, (-node_addr - 1));
|
||||
int prim_addr = __float_as_int(leaf.x);
|
||||
|
||||
if (prim_addr >= 0) {
|
||||
@@ -137,7 +137,7 @@ ccl_device_inline
|
||||
|
||||
/* primitive intersection */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert((kernel_tex_fetch(__prim_type, prim_addr) & PRIMITIVE_ALL) ==
|
||||
kernel_assert((kernel_data_fetch(prim_type, prim_addr) & PRIMITIVE_ALL) ==
|
||||
(type & PRIMITIVE_ALL));
|
||||
bool hit;
|
||||
|
||||
@@ -147,9 +147,9 @@ ccl_device_inline
|
||||
Intersection isect ccl_optional_struct_init;
|
||||
|
||||
const int prim_object = (object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, prim_addr) :
|
||||
kernel_data_fetch(prim_object, prim_addr) :
|
||||
object;
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self_shadow(ray->self, prim_object, prim)) {
|
||||
continue;
|
||||
}
|
||||
@@ -181,14 +181,14 @@ ccl_device_inline
|
||||
case PRIMITIVE_CURVE_RIBBON:
|
||||
case PRIMITIVE_MOTION_CURVE_RIBBON: {
|
||||
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
|
||||
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
|
||||
const float2 prim_time = kernel_data_fetch(prim_time, prim_addr);
|
||||
if (ray->time < prim_time.x || ray->time > prim_time.y) {
|
||||
hit = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
|
||||
const int curve_type = kernel_data_fetch(prim_type, prim_addr);
|
||||
hit = curve_intersect(
|
||||
kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, curve_type);
|
||||
|
||||
@@ -199,14 +199,14 @@ ccl_device_inline
|
||||
case PRIMITIVE_POINT:
|
||||
case PRIMITIVE_MOTION_POINT: {
|
||||
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
|
||||
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
|
||||
const float2 prim_time = kernel_data_fetch(prim_time, prim_addr);
|
||||
if (ray->time < prim_time.x || ray->time > prim_time.y) {
|
||||
hit = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
|
||||
const int point_type = kernel_data_fetch(prim_type, prim_addr);
|
||||
hit = point_intersect(
|
||||
kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, point_type);
|
||||
break;
|
||||
@@ -291,7 +291,7 @@ ccl_device_inline
|
||||
}
|
||||
else {
|
||||
/* instance push */
|
||||
object = kernel_tex_fetch(__prim_object, -prim_addr - 1);
|
||||
object = kernel_data_fetch(prim_object, -prim_addr - 1);
|
||||
|
||||
#if BVH_FEATURE(BVH_MOTION)
|
||||
t_world_to_instance = bvh_instance_motion_push(
|
||||
@@ -307,7 +307,7 @@ ccl_device_inline
|
||||
kernel_assert(stack_ptr < BVH_STACK_SIZE);
|
||||
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
|
||||
|
||||
node_addr = kernel_tex_fetch(__object_node, object);
|
||||
node_addr = kernel_data_fetch(object_node, object);
|
||||
}
|
||||
}
|
||||
} while (node_addr != ENTRYPOINT_SENTINEL);
|
||||
|
@@ -62,7 +62,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
while (node_addr >= 0 && node_addr != ENTRYPOINT_SENTINEL) {
|
||||
int node_addr_child1, traverse_mask;
|
||||
float dist[2];
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
|
||||
{
|
||||
traverse_mask = NODE_INTERSECT(kg,
|
||||
@@ -108,7 +108,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
|
||||
/* if node is leaf, fetch triangle list */
|
||||
if (node_addr < 0) {
|
||||
float4 leaf = kernel_tex_fetch(__bvh_leaf_nodes, (-node_addr - 1));
|
||||
float4 leaf = kernel_data_fetch(bvh_leaf_nodes, (-node_addr - 1));
|
||||
int prim_addr = __float_as_int(leaf.x);
|
||||
|
||||
if (prim_addr >= 0) {
|
||||
@@ -121,12 +121,12 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
|
||||
/* primitive intersection */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
|
||||
const int prim_object = (object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, prim_addr) :
|
||||
kernel_data_fetch(prim_object, prim_addr) :
|
||||
object;
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self_shadow(ray->self, prim_object, prim)) {
|
||||
continue;
|
||||
}
|
||||
@@ -166,13 +166,13 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
case PRIMITIVE_CURVE_RIBBON:
|
||||
case PRIMITIVE_MOTION_CURVE_RIBBON: {
|
||||
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
|
||||
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
|
||||
const float2 prim_time = kernel_data_fetch(prim_time, prim_addr);
|
||||
if (ray->time < prim_time.x || ray->time > prim_time.y) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
const int curve_type = kernel_tex_fetch(__prim_type, prim_addr);
|
||||
const int curve_type = kernel_data_fetch(prim_type, prim_addr);
|
||||
const bool hit = curve_intersect(
|
||||
kg, isect, P, dir, isect->t, prim_object, prim, ray->time, curve_type);
|
||||
if (hit) {
|
||||
@@ -187,13 +187,13 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
case PRIMITIVE_POINT:
|
||||
case PRIMITIVE_MOTION_POINT: {
|
||||
if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) {
|
||||
const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr);
|
||||
const float2 prim_time = kernel_data_fetch(prim_time, prim_addr);
|
||||
if (ray->time < prim_time.x || ray->time > prim_time.y) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
const int point_type = kernel_tex_fetch(__prim_type, prim_addr);
|
||||
const int point_type = kernel_data_fetch(prim_type, prim_addr);
|
||||
const bool hit = point_intersect(
|
||||
kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type);
|
||||
if (hit) {
|
||||
@@ -209,7 +209,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
}
|
||||
else {
|
||||
/* instance push */
|
||||
object = kernel_tex_fetch(__prim_object, -prim_addr - 1);
|
||||
object = kernel_data_fetch(prim_object, -prim_addr - 1);
|
||||
|
||||
#if BVH_FEATURE(BVH_MOTION)
|
||||
isect->t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm);
|
||||
@@ -221,7 +221,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
|
||||
kernel_assert(stack_ptr < BVH_STACK_SIZE);
|
||||
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
|
||||
|
||||
node_addr = kernel_tex_fetch(__object_node, object);
|
||||
node_addr = kernel_data_fetch(object_node, object);
|
||||
}
|
||||
}
|
||||
} while (node_addr != ENTRYPOINT_SENTINEL);
|
||||
|
@@ -53,20 +53,20 @@ ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg,
|
||||
int shader = 0;
|
||||
|
||||
if (type & PRIMITIVE_TRIANGLE) {
|
||||
shader = kernel_tex_fetch(__tri_shader, prim);
|
||||
shader = kernel_data_fetch(tri_shader, prim);
|
||||
}
|
||||
#ifdef __POINTCLOUD__
|
||||
else if (type & PRIMITIVE_POINT) {
|
||||
shader = kernel_tex_fetch(__points_shader, prim);
|
||||
shader = kernel_data_fetch(points_shader, prim);
|
||||
}
|
||||
#endif
|
||||
#ifdef __HAIR__
|
||||
else if (type & PRIMITIVE_CURVE) {
|
||||
shader = kernel_tex_fetch(__curves, prim).shader_id;
|
||||
shader = kernel_data_fetch(curves, prim).shader_id;
|
||||
}
|
||||
#endif
|
||||
|
||||
return kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).flags;
|
||||
return kernel_data_fetch(shaders, (shader & SHADER_MASK)).flags;
|
||||
}
|
||||
|
||||
ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals kg,
|
||||
@@ -76,16 +76,16 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals
|
||||
int shader = 0;
|
||||
|
||||
if (isect_type & PRIMITIVE_TRIANGLE) {
|
||||
shader = kernel_tex_fetch(__tri_shader, prim);
|
||||
shader = kernel_data_fetch(tri_shader, prim);
|
||||
}
|
||||
#ifdef __POINTCLOUD__
|
||||
else if (isect_type & PRIMITIVE_POINT) {
|
||||
shader = kernel_tex_fetch(__points_shader, prim);
|
||||
shader = kernel_data_fetch(points_shader, prim);
|
||||
}
|
||||
#endif
|
||||
#ifdef __HAIR__
|
||||
else if (isect_type & PRIMITIVE_CURVE) {
|
||||
shader = kernel_tex_fetch(__curves, prim).shader_id;
|
||||
shader = kernel_data_fetch(curves, prim).shader_id;
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -101,7 +101,7 @@ ccl_device_forceinline int intersection_get_shader(
|
||||
ccl_device_forceinline int intersection_get_object_flags(
|
||||
KernelGlobals kg, ccl_private const Intersection *ccl_restrict isect)
|
||||
{
|
||||
return kernel_tex_fetch(__object_flag, isect->object);
|
||||
return kernel_data_fetch(object_flag, isect->object);
|
||||
}
|
||||
|
||||
/* TODO: find a better (faster) solution for this. Maybe store offset per object for
|
||||
@@ -110,8 +110,8 @@ ccl_device_inline int intersection_find_attribute(KernelGlobals kg,
|
||||
const int object,
|
||||
const uint id)
|
||||
{
|
||||
uint attr_offset = kernel_tex_fetch(__objects, object).attribute_map_offset;
|
||||
AttributeMap attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
|
||||
uint attr_offset = kernel_data_fetch(objects, object).attribute_map_offset;
|
||||
AttributeMap attr_map = kernel_data_fetch(attributes_map, attr_offset);
|
||||
|
||||
while (attr_map.id != id) {
|
||||
if (UNLIKELY(attr_map.id == ATTR_STD_NONE)) {
|
||||
@@ -126,7 +126,7 @@ ccl_device_inline int intersection_find_attribute(KernelGlobals kg,
|
||||
else {
|
||||
attr_offset += ATTR_PRIM_TYPES;
|
||||
}
|
||||
attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
|
||||
attr_map = kernel_data_fetch(attributes_map, attr_offset);
|
||||
}
|
||||
|
||||
/* return result */
|
||||
@@ -151,12 +151,12 @@ ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg,
|
||||
}
|
||||
|
||||
/* Interpolate transparency between curve keys. */
|
||||
const KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
|
||||
const KernelCurve kcurve = kernel_data_fetch(curves, prim);
|
||||
const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(kcurve.type);
|
||||
const int k1 = k0 + 1;
|
||||
|
||||
const float f0 = kernel_tex_fetch(__attributes_float, offset + k0);
|
||||
const float f1 = kernel_tex_fetch(__attributes_float, offset + k1);
|
||||
const float f0 = kernel_data_fetch(attributes_float, offset + k0);
|
||||
const float f1 = kernel_data_fetch(attributes_float, offset + k1);
|
||||
|
||||
return (1.0f - u) * f0 + u * f1;
|
||||
}
|
||||
|
@@ -65,7 +65,7 @@ ccl_device_inline
|
||||
while (node_addr >= 0 && node_addr != ENTRYPOINT_SENTINEL) {
|
||||
int node_addr_child1, traverse_mask;
|
||||
float dist[2];
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
|
||||
traverse_mask = NODE_INTERSECT(kg,
|
||||
P,
|
||||
@@ -109,7 +109,7 @@ ccl_device_inline
|
||||
|
||||
/* if node is leaf, fetch triangle list */
|
||||
if (node_addr < 0) {
|
||||
float4 leaf = kernel_tex_fetch(__bvh_leaf_nodes, (-node_addr - 1));
|
||||
float4 leaf = kernel_data_fetch(bvh_leaf_nodes, (-node_addr - 1));
|
||||
int prim_addr = __float_as_int(leaf.x);
|
||||
|
||||
if (prim_addr >= 0) {
|
||||
@@ -125,17 +125,17 @@ ccl_device_inline
|
||||
case PRIMITIVE_TRIANGLE: {
|
||||
/* intersect ray against primitive */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
/* only primitives from volume object */
|
||||
const int prim_object = (object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, prim_addr) :
|
||||
kernel_data_fetch(prim_object, prim_addr) :
|
||||
object;
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self(ray->self, prim_object, prim)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
|
||||
int object_flag = kernel_data_fetch(object_flag, prim_object);
|
||||
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
continue;
|
||||
}
|
||||
@@ -148,16 +148,16 @@ ccl_device_inline
|
||||
case PRIMITIVE_MOTION_TRIANGLE: {
|
||||
/* intersect ray against primitive */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
/* only primitives from volume object */
|
||||
const int prim_object = (object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, prim_addr) :
|
||||
kernel_data_fetch(prim_object, prim_addr) :
|
||||
object;
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self(ray->self, prim_object, prim)) {
|
||||
continue;
|
||||
}
|
||||
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
|
||||
int object_flag = kernel_data_fetch(object_flag, prim_object);
|
||||
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
continue;
|
||||
}
|
||||
@@ -182,8 +182,8 @@ ccl_device_inline
|
||||
}
|
||||
else {
|
||||
/* instance push */
|
||||
object = kernel_tex_fetch(__prim_object, -prim_addr - 1);
|
||||
int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
object = kernel_data_fetch(prim_object, -prim_addr - 1);
|
||||
int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (object_flag & SD_OBJECT_HAS_VOLUME) {
|
||||
#if BVH_FEATURE(BVH_MOTION)
|
||||
isect->t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm);
|
||||
@@ -195,7 +195,7 @@ ccl_device_inline
|
||||
kernel_assert(stack_ptr < BVH_STACK_SIZE);
|
||||
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
|
||||
|
||||
node_addr = kernel_tex_fetch(__object_node, object);
|
||||
node_addr = kernel_data_fetch(object_node, object);
|
||||
}
|
||||
else {
|
||||
/* pop */
|
||||
|
@@ -67,7 +67,7 @@ ccl_device_inline
|
||||
while (node_addr >= 0 && node_addr != ENTRYPOINT_SENTINEL) {
|
||||
int node_addr_child1, traverse_mask;
|
||||
float dist[2];
|
||||
float4 cnodes = kernel_tex_fetch(__bvh_nodes, node_addr + 0);
|
||||
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
|
||||
|
||||
traverse_mask = NODE_INTERSECT(kg,
|
||||
P,
|
||||
@@ -111,7 +111,7 @@ ccl_device_inline
|
||||
|
||||
/* if node is leaf, fetch triangle list */
|
||||
if (node_addr < 0) {
|
||||
float4 leaf = kernel_tex_fetch(__bvh_leaf_nodes, (-node_addr - 1));
|
||||
float4 leaf = kernel_data_fetch(bvh_leaf_nodes, (-node_addr - 1));
|
||||
int prim_addr = __float_as_int(leaf.x);
|
||||
|
||||
if (prim_addr >= 0) {
|
||||
@@ -128,16 +128,16 @@ ccl_device_inline
|
||||
case PRIMITIVE_TRIANGLE: {
|
||||
/* intersect ray against primitive */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
/* only primitives from volume object */
|
||||
const int prim_object = (object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, prim_addr) :
|
||||
kernel_data_fetch(prim_object, prim_addr) :
|
||||
object;
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self(ray->self, prim_object, prim)) {
|
||||
continue;
|
||||
}
|
||||
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
|
||||
int object_flag = kernel_data_fetch(object_flag, prim_object);
|
||||
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
continue;
|
||||
}
|
||||
@@ -172,16 +172,16 @@ ccl_device_inline
|
||||
case PRIMITIVE_MOTION_TRIANGLE: {
|
||||
/* intersect ray against primitive */
|
||||
for (; prim_addr < prim_addr2; prim_addr++) {
|
||||
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
|
||||
kernel_assert(kernel_data_fetch(prim_type, prim_addr) == type);
|
||||
/* only primitives from volume object */
|
||||
const int prim_object = (object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, prim_addr) :
|
||||
kernel_data_fetch(prim_object, prim_addr) :
|
||||
object;
|
||||
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
|
||||
const int prim = kernel_data_fetch(prim_index, prim_addr);
|
||||
if (intersection_skip_self(ray->self, prim_object, prim)) {
|
||||
continue;
|
||||
}
|
||||
int object_flag = kernel_tex_fetch(__object_flag, prim_object);
|
||||
int object_flag = kernel_data_fetch(object_flag, prim_object);
|
||||
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
continue;
|
||||
}
|
||||
@@ -228,8 +228,8 @@ ccl_device_inline
|
||||
}
|
||||
else {
|
||||
/* instance push */
|
||||
object = kernel_tex_fetch(__prim_object, -prim_addr - 1);
|
||||
int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
object = kernel_data_fetch(prim_object, -prim_addr - 1);
|
||||
int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (object_flag & SD_OBJECT_HAS_VOLUME) {
|
||||
#if BVH_FEATURE(BVH_MOTION)
|
||||
isect_t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm);
|
||||
@@ -244,7 +244,7 @@ ccl_device_inline
|
||||
kernel_assert(stack_ptr < BVH_STACK_SIZE);
|
||||
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
|
||||
|
||||
node_addr = kernel_tex_fetch(__object_node, object);
|
||||
node_addr = kernel_data_fetch(object_node, object);
|
||||
}
|
||||
else {
|
||||
/* pop */
|
||||
|
@@ -90,7 +90,7 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
|
||||
#ifdef __CAMERA_MOTION__
|
||||
if (kernel_data.cam.num_motion_steps) {
|
||||
transform_motion_array_interpolate(&cameratoworld,
|
||||
kernel_tex_array(__camera_motion),
|
||||
kernel_data_array(camera_motion),
|
||||
kernel_data.cam.num_motion_steps,
|
||||
ray->time);
|
||||
}
|
||||
@@ -210,7 +210,7 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
|
||||
#ifdef __CAMERA_MOTION__
|
||||
if (kernel_data.cam.num_motion_steps) {
|
||||
transform_motion_array_interpolate(&cameratoworld,
|
||||
kernel_tex_array(__camera_motion),
|
||||
kernel_data_array(camera_motion),
|
||||
kernel_data.cam.num_motion_steps,
|
||||
ray->time);
|
||||
}
|
||||
@@ -421,7 +421,7 @@ ccl_device_inline void camera_sample(KernelGlobals kg,
|
||||
}
|
||||
else {
|
||||
#ifdef __CAMERA_MOTION__
|
||||
ccl_global const DecomposedTransform *cam_motion = kernel_tex_array(__camera_motion);
|
||||
ccl_global const DecomposedTransform *cam_motion = kernel_data_array(camera_motion);
|
||||
camera_sample_panorama(&kernel_data.cam, cam_motion, raster_x, raster_y, lens_u, lens_v, ray);
|
||||
#else
|
||||
camera_sample_panorama(&kernel_data.cam, raster_x, raster_y, lens_u, lens_v, ray);
|
||||
|
@@ -434,7 +434,7 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
|
||||
else {
|
||||
/* Shadow terminator offset. */
|
||||
const float frequency_multiplier =
|
||||
kernel_tex_fetch(__objects, sd->object).shadow_terminator_shading_offset;
|
||||
kernel_data_fetch(objects, sd->object).shadow_terminator_shading_offset;
|
||||
if (frequency_multiplier > 1.0f) {
|
||||
*eval *= shift_cos_in(dot(*omega_in, sc->N), frequency_multiplier);
|
||||
}
|
||||
@@ -556,7 +556,7 @@ ccl_device_inline
|
||||
}
|
||||
/* Shadow terminator offset. */
|
||||
const float frequency_multiplier =
|
||||
kernel_tex_fetch(__objects, sd->object).shadow_terminator_shading_offset;
|
||||
kernel_data_fetch(objects, sd->object).shadow_terminator_shading_offset;
|
||||
if (frequency_multiplier > 1.0f) {
|
||||
eval *= shift_cos_in(dot(omega_in, sc->N), frequency_multiplier);
|
||||
}
|
||||
|
82
intern/cycles/kernel/data_arrays.h
Normal file
82
intern/cycles/kernel/data_arrays.h
Normal file
@@ -0,0 +1,82 @@
|
||||
/* SPDX-License-Identifier: Apache-2.0
|
||||
* Copyright 2011-2022 Blender Foundation */
|
||||
|
||||
#ifndef KERNEL_DATA_ARRAY
|
||||
# define KERNEL_DATA_ARRAY(type, name)
|
||||
#endif
|
||||
|
||||
/* BVH2, not used for OptiX or Embree. */
|
||||
KERNEL_DATA_ARRAY(float4, bvh_nodes)
|
||||
KERNEL_DATA_ARRAY(float4, bvh_leaf_nodes)
|
||||
KERNEL_DATA_ARRAY(uint, prim_type)
|
||||
KERNEL_DATA_ARRAY(uint, prim_visibility)
|
||||
KERNEL_DATA_ARRAY(uint, prim_index)
|
||||
KERNEL_DATA_ARRAY(uint, prim_object)
|
||||
KERNEL_DATA_ARRAY(uint, object_node)
|
||||
KERNEL_DATA_ARRAY(float2, prim_time)
|
||||
|
||||
/* objects */
|
||||
KERNEL_DATA_ARRAY(KernelObject, objects)
|
||||
KERNEL_DATA_ARRAY(Transform, object_motion_pass)
|
||||
KERNEL_DATA_ARRAY(DecomposedTransform, object_motion)
|
||||
KERNEL_DATA_ARRAY(uint, object_flag)
|
||||
KERNEL_DATA_ARRAY(float, object_volume_step)
|
||||
KERNEL_DATA_ARRAY(uint, object_prim_offset)
|
||||
|
||||
/* cameras */
|
||||
KERNEL_DATA_ARRAY(DecomposedTransform, camera_motion)
|
||||
|
||||
/* triangles */
|
||||
KERNEL_DATA_ARRAY(uint, tri_shader)
|
||||
KERNEL_DATA_ARRAY(packed_float3, tri_vnormal)
|
||||
KERNEL_DATA_ARRAY(uint4, tri_vindex)
|
||||
KERNEL_DATA_ARRAY(uint, tri_patch)
|
||||
KERNEL_DATA_ARRAY(float2, tri_patch_uv)
|
||||
KERNEL_DATA_ARRAY(packed_float3, tri_verts)
|
||||
|
||||
/* curves */
|
||||
KERNEL_DATA_ARRAY(KernelCurve, curves)
|
||||
KERNEL_DATA_ARRAY(float4, curve_keys)
|
||||
KERNEL_DATA_ARRAY(KernelCurveSegment, curve_segments)
|
||||
|
||||
/* patches */
|
||||
KERNEL_DATA_ARRAY(uint, patches)
|
||||
|
||||
/* pointclouds */
|
||||
KERNEL_DATA_ARRAY(float4, points)
|
||||
KERNEL_DATA_ARRAY(uint, points_shader)
|
||||
|
||||
/* attributes */
|
||||
KERNEL_DATA_ARRAY(AttributeMap, attributes_map)
|
||||
KERNEL_DATA_ARRAY(float, attributes_float)
|
||||
KERNEL_DATA_ARRAY(float2, attributes_float2)
|
||||
KERNEL_DATA_ARRAY(packed_float3, attributes_float3)
|
||||
KERNEL_DATA_ARRAY(float4, attributes_float4)
|
||||
KERNEL_DATA_ARRAY(uchar4, attributes_uchar4)
|
||||
|
||||
/* lights */
|
||||
KERNEL_DATA_ARRAY(KernelLightDistribution, light_distribution)
|
||||
KERNEL_DATA_ARRAY(KernelLight, lights)
|
||||
KERNEL_DATA_ARRAY(float2, light_background_marginal_cdf)
|
||||
KERNEL_DATA_ARRAY(float2, light_background_conditional_cdf)
|
||||
|
||||
/* particles */
|
||||
KERNEL_DATA_ARRAY(KernelParticle, particles)
|
||||
|
||||
/* shaders */
|
||||
KERNEL_DATA_ARRAY(uint4, svm_nodes)
|
||||
KERNEL_DATA_ARRAY(KernelShader, shaders)
|
||||
|
||||
/* lookup tables */
|
||||
KERNEL_DATA_ARRAY(float, lookup_table)
|
||||
|
||||
/* sobol */
|
||||
KERNEL_DATA_ARRAY(float, sample_pattern_lut)
|
||||
|
||||
/* image textures */
|
||||
KERNEL_DATA_ARRAY(TextureInfo, texture_info)
|
||||
|
||||
/* ies lights */
|
||||
KERNEL_DATA_ARRAY(float, ies)
|
||||
|
||||
#undef KERNEL_DATA_ARRAY
|
@@ -35,20 +35,6 @@ CCL_NAMESPACE_BEGIN
|
||||
|
||||
#define kernel_assert(cond) assert(cond)
|
||||
|
||||
/* Texture types to be compatible with CUDA textures. These are really just
|
||||
* simple arrays and after inlining fetch hopefully revert to being a simple
|
||||
* pointer lookup. */
|
||||
template<typename T> struct texture {
|
||||
ccl_always_inline const T &fetch(int index) const
|
||||
{
|
||||
kernel_assert(index >= 0 && index < width);
|
||||
return data[index];
|
||||
}
|
||||
|
||||
T *data;
|
||||
int width;
|
||||
};
|
||||
|
||||
/* Macros to handle different memory storage on different devices */
|
||||
|
||||
#ifdef __KERNEL_SSE2__
|
||||
|
@@ -12,7 +12,7 @@
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
/* On the CPU, we pass along the struct KernelGlobals to nearly everywhere in
|
||||
* the kernel, to access constant data. These are all stored as "textures", but
|
||||
* the kernel, to access constant data. These are all stored as flat arrays.
|
||||
* these are really just standard arrays. We can't use actually globals because
|
||||
* multiple renders may be running inside the same process. */
|
||||
|
||||
@@ -22,11 +22,23 @@ struct OSLThreadData;
|
||||
struct OSLShadingSystem;
|
||||
#endif
|
||||
|
||||
typedef struct KernelGlobalsCPU {
|
||||
#define KERNEL_TEX(type, name) texture<type> name;
|
||||
#include "kernel/textures.h"
|
||||
/* Array for kernel data, with size to be able to assert on invalid data access. */
|
||||
template<typename T> struct kernel_array {
|
||||
ccl_always_inline const T &fetch(int index) const
|
||||
{
|
||||
kernel_assert(index >= 0 && index < width);
|
||||
return data[index];
|
||||
}
|
||||
|
||||
KernelData __data;
|
||||
T *data;
|
||||
int width;
|
||||
};
|
||||
|
||||
typedef struct KernelGlobalsCPU {
|
||||
#define KERNEL_DATA_ARRAY(type, name) kernel_array<type> name;
|
||||
#include "kernel/data_arrays.h"
|
||||
|
||||
KernelData data;
|
||||
|
||||
#ifdef __OSL__
|
||||
/* On the CPU, we also have the OSL globals here. Most data structures are shared
|
||||
@@ -44,8 +56,8 @@ typedef struct KernelGlobalsCPU {
|
||||
typedef const KernelGlobalsCPU *ccl_restrict KernelGlobals;
|
||||
|
||||
/* Abstraction macros */
|
||||
#define kernel_tex_fetch(tex, index) (kg->tex.fetch(index))
|
||||
#define kernel_tex_array(tex) (kg->tex.data)
|
||||
#define kernel_data (kg->__data)
|
||||
#define kernel_data_fetch(name, index) (kg->name.fetch(index))
|
||||
#define kernel_data_array(name) (kg->name.data)
|
||||
#define kernel_data (kg->data)
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -733,7 +733,7 @@ template<typename TexT, typename OutT = float4> struct NanoVDBInterpolator {
|
||||
|
||||
ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, int id, float x, float y)
|
||||
{
|
||||
const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
const TextureInfo &info = kernel_data_fetch(texture_info, id);
|
||||
|
||||
if (UNLIKELY(!info.data)) {
|
||||
return zero_float4();
|
||||
@@ -776,7 +776,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
|
||||
float3 P,
|
||||
InterpolationType interp)
|
||||
{
|
||||
const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
const TextureInfo &info = kernel_data_fetch(texture_info, id);
|
||||
|
||||
if (UNLIKELY(!info.data)) {
|
||||
return zero_float4();
|
||||
|
@@ -53,8 +53,8 @@ CCL_NAMESPACE_BEGIN
|
||||
|
||||
void kernel_const_copy(KernelGlobalsCPU *kg, const char *name, void *host, size_t)
|
||||
{
|
||||
if (strcmp(name, "__data") == 0) {
|
||||
kg->__data = *(KernelData *)host;
|
||||
if (strcmp(name, "data") == 0) {
|
||||
kg->data = *(KernelData *)host;
|
||||
}
|
||||
else {
|
||||
assert(0);
|
||||
@@ -66,13 +66,13 @@ void kernel_global_memory_copy(KernelGlobalsCPU *kg, const char *name, void *mem
|
||||
if (0) {
|
||||
}
|
||||
|
||||
#define KERNEL_TEX(type, tname) \
|
||||
#define KERNEL_DATA_ARRAY(type, tname) \
|
||||
else if (strcmp(name, #tname) == 0) \
|
||||
{ \
|
||||
kg->tname.data = (type *)mem; \
|
||||
kg->tname.width = size; \
|
||||
}
|
||||
#include "kernel/textures.h"
|
||||
#include "kernel/data_arrays.h"
|
||||
else {
|
||||
assert(0);
|
||||
}
|
||||
|
@@ -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 */
|
||||
__constant__ IntegratorStateGPU __integrator_state;
|
||||
/* Integrator state */
|
||||
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
|
||||
|
@@ -181,7 +181,7 @@ ccl_device_noinline typename nanovdb::NanoGrid<T>::ValueType kernel_tex_image_in
|
||||
|
||||
ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, int id, float x, float y)
|
||||
{
|
||||
ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
ccl_global const TextureInfo &info = kernel_data_fetch(texture_info, id);
|
||||
|
||||
/* float4, byte4, ushort4 and half4 */
|
||||
const int texture_type = info.data_type;
|
||||
@@ -216,7 +216,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
|
||||
float3 P,
|
||||
InterpolationType interp)
|
||||
{
|
||||
ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
|
||||
ccl_global const TextureInfo &info = kernel_data_fetch(texture_info, id);
|
||||
|
||||
if (info.use_transform_3d) {
|
||||
P = transform_point(&info.transform_3d, P);
|
||||
|
@@ -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 */
|
||||
__constant__ IntegratorStateGPU __integrator_state;
|
||||
/* Integrator state */
|
||||
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
|
||||
|
@@ -7,4 +7,4 @@
|
||||
/* NOTE: These macros will need maintaining as entry-points change. */
|
||||
|
||||
#undef kernel_integrator_state
|
||||
#define kernel_integrator_state context.launch_params_metal.__integrator_state
|
||||
#define kernel_integrator_state context.launch_params_metal.integrator_state
|
||||
|
@@ -12,11 +12,11 @@ CCL_NAMESPACE_BEGIN
|
||||
|
||||
typedef struct KernelParamsMetal {
|
||||
|
||||
#define KERNEL_TEX(type, name) ccl_global const type *name;
|
||||
#include "kernel/textures.h"
|
||||
#undef KERNEL_TEX
|
||||
#define KERNEL_DATA_ARRAY(type, name) ccl_global const type *name;
|
||||
#include "kernel/data_arrays.h"
|
||||
#undef KERNEL_DATA_ARRAY
|
||||
|
||||
const IntegratorStateGPU __integrator_state;
|
||||
const IntegratorStateGPU integrator_state;
|
||||
const KernelData data;
|
||||
|
||||
} KernelParamsMetal;
|
||||
@@ -27,12 +27,10 @@ typedef struct KernelGlobalsGPU {
|
||||
|
||||
typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
|
||||
|
||||
/* Abstraction macros */
|
||||
#define kernel_data launch_params_metal.data
|
||||
#define kernel_integrator_state launch_params_metal.__integrator_state
|
||||
|
||||
/* data lookup defines */
|
||||
|
||||
#define kernel_tex_fetch(tex, index) launch_params_metal.tex[index]
|
||||
#define kernel_tex_array(tex) launch_params_metal.tex
|
||||
#define kernel_data_fetch(name, index) launch_params_metal.name[index]
|
||||
#define kernel_data_array(name) launch_params_metal.name
|
||||
#define kernel_integrator_state launch_params_metal.integrator_state
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -59,7 +59,7 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
|
||||
TReturn result;
|
||||
|
||||
#ifdef __BVH_LOCAL__
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
|
||||
if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) {
|
||||
/* Only intersect with matching object and skip self-intersecton. */
|
||||
@@ -113,16 +113,16 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
|
||||
isect->t = ray_tmax;
|
||||
isect->prim = prim;
|
||||
isect->object = object;
|
||||
isect->type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
isect->type = kernel_data_fetch(objects, object).primitive_type;
|
||||
|
||||
isect->u = 1.0f - barycentrics.y - barycentrics.x;
|
||||
isect->v = barycentrics.x;
|
||||
|
||||
/* Record geometric normal */
|
||||
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect->prim).w;
|
||||
const float3 tri_a = float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0));
|
||||
const float3 tri_b = float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1));
|
||||
const float3 tri_c = float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
|
||||
const uint tri_vindex = kernel_data_fetch(tri_vindex, isect->prim).w;
|
||||
const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex + 0));
|
||||
const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex + 1));
|
||||
const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex + 2));
|
||||
payload.local_isect.Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
|
||||
|
||||
/* Continue tracing (without this the trace call would return after the first hit) */
|
||||
@@ -168,7 +168,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
|
||||
#ifdef __SHADOW_RECORD_ALL__
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = payload.visibility;
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
/* continue search */
|
||||
return true;
|
||||
}
|
||||
@@ -184,14 +184,14 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
|
||||
if (intersection_type == METALRT_HIT_TRIANGLE) {
|
||||
u = 1.0f - barycentrics.y - barycentrics.x;
|
||||
v = barycentrics.x;
|
||||
type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
type = kernel_data_fetch(objects, object).primitive_type;
|
||||
}
|
||||
# ifdef __HAIR__
|
||||
else {
|
||||
u = barycentrics.x;
|
||||
v = barycentrics.y;
|
||||
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
type = segment.type;
|
||||
prim = segment.prim;
|
||||
|
||||
@@ -294,7 +294,7 @@ __anyhit__cycles_metalrt_shadow_all_hit_tri(constant KernelParamsMetal &launch_p
|
||||
float2 barycentrics [[barycentric_coord]],
|
||||
float ray_tmax [[distance]])
|
||||
{
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
|
||||
TriangleIntersectionResult result;
|
||||
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_TRIANGLE>(
|
||||
@@ -337,7 +337,7 @@ inline TReturnType metalrt_visibility_test(constant KernelParamsMetal &launch_pa
|
||||
|
||||
uint visibility = payload.visibility;
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
result.accept = false;
|
||||
result.continue_search = true;
|
||||
return result;
|
||||
@@ -377,12 +377,12 @@ __anyhit__cycles_metalrt_visibility_test_tri(constant KernelParamsMetal &launch_
|
||||
unsigned int object [[user_instance_id]],
|
||||
unsigned int primitive_id [[primitive_id]])
|
||||
{
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
TriangleIntersectionResult result = metalrt_visibility_test<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
|
||||
launch_params_metal, payload, object, prim, 0.0f);
|
||||
if (result.accept) {
|
||||
payload.prim = prim;
|
||||
payload.type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
payload.type = kernel_data_fetch(objects, object).primitive_type;
|
||||
}
|
||||
return result;
|
||||
}
|
||||
@@ -414,7 +414,7 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal,
|
||||
{
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = payload.visibility;
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return;
|
||||
}
|
||||
# endif
|
||||
@@ -495,8 +495,8 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
@@ -526,8 +526,8 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
@@ -557,8 +557,8 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
@@ -585,8 +585,8 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
@@ -620,7 +620,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
|
||||
{
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = payload.visibility;
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return;
|
||||
}
|
||||
# endif
|
||||
@@ -701,8 +701,8 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
const uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const int type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
const int type = kernel_data_fetch(objects, object).primitive_type;
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
@@ -730,8 +730,8 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
|
||||
const float3 ray_direction [[direction]],
|
||||
const float ray_tmax [[max_distance]])
|
||||
{
|
||||
const uint prim = primitive_id + kernel_tex_fetch(__object_prim_offset, object);
|
||||
const int type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
|
||||
const int type = kernel_data_fetch(objects, object).primitive_type;
|
||||
|
||||
BoundingBoxIntersectionResult result;
|
||||
result.accept = false;
|
||||
|
@@ -28,21 +28,21 @@ struct KernelParamsOptiX {
|
||||
|
||||
/* Global scene data and textures */
|
||||
KernelData data;
|
||||
#define KERNEL_TEX(type, name) const type *name;
|
||||
#include "kernel/textures.h"
|
||||
#define KERNEL_DATA_ARRAY(type, name) const type *name;
|
||||
#include "kernel/data_arrays.h"
|
||||
|
||||
/* Integrator state */
|
||||
IntegratorStateGPU __integrator_state;
|
||||
IntegratorStateGPU integrator_state;
|
||||
};
|
||||
|
||||
#ifdef __NVCC__
|
||||
extern "C" static __constant__ KernelParamsOptiX __params;
|
||||
extern "C" static __constant__ KernelParamsOptiX kernel_params;
|
||||
#endif
|
||||
|
||||
/* Abstraction macros */
|
||||
#define kernel_data __params.data
|
||||
#define kernel_tex_array(t) __params.t
|
||||
#define kernel_tex_fetch(t, index) __params.t[(index)]
|
||||
#define kernel_integrator_state __params.__integrator_state
|
||||
#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_integrator_state kernel_params.integrator_state
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -51,15 +51,15 @@ ccl_device_forceinline int get_object_id()
|
||||
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
|
||||
{
|
||||
const int global_index = optixGetLaunchIndex().x;
|
||||
const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
|
||||
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
|
||||
global_index;
|
||||
integrator_intersect_closest(nullptr, path_index, __params.render_buffer);
|
||||
integrator_intersect_closest(nullptr, path_index, kernel_params.render_buffer);
|
||||
}
|
||||
|
||||
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()
|
||||
{
|
||||
const int global_index = optixGetLaunchIndex().x;
|
||||
const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
|
||||
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
|
||||
global_index;
|
||||
integrator_intersect_shadow(nullptr, path_index);
|
||||
}
|
||||
@@ -67,7 +67,7 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()
|
||||
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurface()
|
||||
{
|
||||
const int global_index = optixGetLaunchIndex().x;
|
||||
const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
|
||||
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
|
||||
global_index;
|
||||
integrator_intersect_subsurface(nullptr, path_index);
|
||||
}
|
||||
@@ -75,7 +75,7 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurfac
|
||||
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_stack()
|
||||
{
|
||||
const int global_index = optixGetLaunchIndex().x;
|
||||
const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
|
||||
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
|
||||
global_index;
|
||||
integrator_intersect_volume_stack(nullptr, path_index);
|
||||
}
|
||||
@@ -151,17 +151,17 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
|
||||
isect->t = optixGetRayTmax();
|
||||
isect->prim = prim;
|
||||
isect->object = get_object_id();
|
||||
isect->type = kernel_tex_fetch(__objects, isect->object).primitive_type;
|
||||
isect->type = kernel_data_fetch(objects, isect->object).primitive_type;
|
||||
|
||||
const float2 barycentrics = optixGetTriangleBarycentrics();
|
||||
isect->u = 1.0f - barycentrics.y - barycentrics.x;
|
||||
isect->v = barycentrics.x;
|
||||
|
||||
/* Record geometric normal. */
|
||||
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
|
||||
const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0);
|
||||
const float3 tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1);
|
||||
const float3 tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
|
||||
const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w;
|
||||
const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0);
|
||||
const float3 tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1);
|
||||
const float3 tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
|
||||
local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
|
||||
|
||||
/* Continue tracing (without this the trace call would return after the first hit). */
|
||||
@@ -176,7 +176,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
|
||||
const uint object = get_object_id();
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = optixGetPayload_4();
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return optixIgnoreIntersection();
|
||||
}
|
||||
# endif
|
||||
@@ -192,14 +192,14 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
|
||||
const float2 barycentrics = optixGetTriangleBarycentrics();
|
||||
u = 1.0f - barycentrics.y - barycentrics.x;
|
||||
v = barycentrics.x;
|
||||
type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
type = kernel_data_fetch(objects, object).primitive_type;
|
||||
}
|
||||
# ifdef __HAIR__
|
||||
else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
|
||||
u = __uint_as_float(optixGetAttribute_0());
|
||||
v = __uint_as_float(optixGetAttribute_1());
|
||||
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
type = segment.type;
|
||||
prim = segment.prim;
|
||||
|
||||
@@ -212,7 +212,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
|
||||
}
|
||||
# endif
|
||||
else {
|
||||
type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
type = kernel_data_fetch(objects, object).primitive_type;
|
||||
u = 0.0f;
|
||||
v = 0.0f;
|
||||
}
|
||||
@@ -307,12 +307,12 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test()
|
||||
const uint object = get_object_id();
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = optixGetPayload_4();
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return optixIgnoreIntersection();
|
||||
}
|
||||
#endif
|
||||
|
||||
if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
if ((kernel_data_fetch(object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
|
||||
return optixIgnoreIntersection();
|
||||
}
|
||||
|
||||
@@ -340,7 +340,7 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
|
||||
const uint object = get_object_id();
|
||||
const uint visibility = optixGetPayload_4();
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return optixIgnoreIntersection();
|
||||
}
|
||||
#endif
|
||||
@@ -377,10 +377,10 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
|
||||
optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x));
|
||||
optixSetPayload_2(__float_as_uint(barycentrics.x));
|
||||
optixSetPayload_3(prim);
|
||||
optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
|
||||
optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
|
||||
}
|
||||
else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, prim);
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
|
||||
optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
|
||||
optixSetPayload_2(optixGetAttribute_1());
|
||||
optixSetPayload_3(segment.prim);
|
||||
@@ -390,7 +390,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
|
||||
optixSetPayload_1(0);
|
||||
optixSetPayload_2(0);
|
||||
optixSetPayload_3(prim);
|
||||
optixSetPayload_5(kernel_tex_fetch(__objects, object).primitive_type);
|
||||
optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -401,7 +401,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
|
||||
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = optixGetPayload_4();
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return;
|
||||
}
|
||||
# endif
|
||||
@@ -436,7 +436,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
|
||||
|
||||
extern "C" __global__ void __intersection__curve_ribbon()
|
||||
{
|
||||
const KernelCurveSegment segment = kernel_tex_fetch(__curve_segments, optixGetPrimitiveIndex());
|
||||
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, optixGetPrimitiveIndex());
|
||||
const int prim = segment.prim;
|
||||
const int type = segment.type;
|
||||
if (type & PRIMITIVE_CURVE_RIBBON) {
|
||||
@@ -451,11 +451,11 @@ extern "C" __global__ void __intersection__point()
|
||||
{
|
||||
const int prim = optixGetPrimitiveIndex();
|
||||
const int object = get_object_id();
|
||||
const int type = kernel_tex_fetch(__objects, object).primitive_type;
|
||||
const int type = kernel_data_fetch(objects, object).primitive_type;
|
||||
|
||||
# ifdef __VISIBILITY_FLAG__
|
||||
const uint visibility = optixGetPayload_4();
|
||||
if ((kernel_tex_fetch(__objects, object).visibility & visibility) == 0) {
|
||||
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
|
||||
return;
|
||||
}
|
||||
# endif
|
||||
|
@@ -11,15 +11,15 @@
|
||||
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_raytrace()
|
||||
{
|
||||
const int global_index = optixGetLaunchIndex().x;
|
||||
const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
|
||||
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
|
||||
global_index;
|
||||
integrator_shade_surface_raytrace(nullptr, path_index, __params.render_buffer);
|
||||
integrator_shade_surface_raytrace(nullptr, path_index, kernel_params.render_buffer);
|
||||
}
|
||||
|
||||
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_mnee()
|
||||
{
|
||||
const int global_index = optixGetLaunchIndex().x;
|
||||
const int path_index = (__params.path_index_array) ? __params.path_index_array[global_index] :
|
||||
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] :
|
||||
global_index;
|
||||
integrator_shade_surface_mnee(nullptr, path_index, __params.render_buffer);
|
||||
integrator_shade_surface_mnee(nullptr, path_index, kernel_params.render_buffer);
|
||||
}
|
||||
|
@@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN
|
||||
|
||||
ccl_device_inline uint subd_triangle_patch(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
return (sd->prim != PRIM_NONE) ? kernel_tex_fetch(__tri_patch, sd->prim) : ~0;
|
||||
return (sd->prim != PRIM_NONE) ? kernel_data_fetch(tri_patch, sd->prim) : ~0;
|
||||
}
|
||||
|
||||
ccl_device_inline uint attribute_primitive_type(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
@@ -42,7 +42,7 @@ ccl_device_inline AttributeDescriptor attribute_not_found()
|
||||
|
||||
ccl_device_inline uint object_attribute_map_offset(KernelGlobals kg, int object)
|
||||
{
|
||||
return kernel_tex_fetch(__objects, object).attribute_map_offset;
|
||||
return kernel_data_fetch(objects, object).attribute_map_offset;
|
||||
}
|
||||
|
||||
ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals kg,
|
||||
@@ -56,7 +56,7 @@ ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals kg,
|
||||
/* for SVM, find attribute by unique id */
|
||||
uint attr_offset = object_attribute_map_offset(kg, sd->object);
|
||||
attr_offset += attribute_primitive_type(kg, sd);
|
||||
AttributeMap attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
|
||||
AttributeMap attr_map = kernel_data_fetch(attributes_map, attr_offset);
|
||||
|
||||
while (attr_map.id != id) {
|
||||
if (UNLIKELY(attr_map.id == ATTR_STD_NONE)) {
|
||||
@@ -71,7 +71,7 @@ ccl_device_inline AttributeDescriptor find_attribute(KernelGlobals kg,
|
||||
else {
|
||||
attr_offset += ATTR_PRIM_TYPES;
|
||||
}
|
||||
attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
|
||||
attr_map = kernel_data_fetch(attributes_map, attr_offset);
|
||||
}
|
||||
|
||||
AttributeDescriptor desc;
|
||||
@@ -99,9 +99,9 @@ ccl_device Transform primitive_attribute_matrix(KernelGlobals kg,
|
||||
{
|
||||
Transform tfm;
|
||||
|
||||
tfm.x = kernel_tex_fetch(__attributes_float4, desc.offset + 0);
|
||||
tfm.y = kernel_tex_fetch(__attributes_float4, desc.offset + 1);
|
||||
tfm.z = kernel_tex_fetch(__attributes_float4, desc.offset + 2);
|
||||
tfm.x = kernel_data_fetch(attributes_float4, desc.offset + 0);
|
||||
tfm.y = kernel_data_fetch(attributes_float4, desc.offset + 1);
|
||||
tfm.z = kernel_data_fetch(attributes_float4, desc.offset + 2);
|
||||
|
||||
return tfm;
|
||||
}
|
||||
|
@@ -23,12 +23,12 @@ ccl_device float curve_attribute_float(KernelGlobals kg,
|
||||
ccl_private float *dy)
|
||||
{
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
|
||||
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
|
||||
KernelCurve curve = kernel_data_fetch(curves, sd->prim);
|
||||
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
|
||||
float f0 = kernel_tex_fetch(__attributes_float, desc.offset + k0);
|
||||
float f1 = kernel_tex_fetch(__attributes_float, desc.offset + k1);
|
||||
float f0 = kernel_data_fetch(attributes_float, desc.offset + k0);
|
||||
float f1 = kernel_data_fetch(attributes_float, desc.offset + k1);
|
||||
|
||||
# ifdef __RAY_DIFFERENTIALS__
|
||||
if (dx)
|
||||
@@ -50,7 +50,7 @@ ccl_device float curve_attribute_float(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float, offset);
|
||||
return kernel_data_fetch(attributes_float, offset);
|
||||
}
|
||||
else {
|
||||
return 0.0f;
|
||||
@@ -65,12 +65,12 @@ ccl_device float2 curve_attribute_float2(KernelGlobals kg,
|
||||
ccl_private float2 *dy)
|
||||
{
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
|
||||
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
|
||||
KernelCurve curve = kernel_data_fetch(curves, sd->prim);
|
||||
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
|
||||
float2 f0 = kernel_tex_fetch(__attributes_float2, desc.offset + k0);
|
||||
float2 f1 = kernel_tex_fetch(__attributes_float2, desc.offset + k1);
|
||||
float2 f0 = kernel_data_fetch(attributes_float2, desc.offset + k0);
|
||||
float2 f1 = kernel_data_fetch(attributes_float2, desc.offset + k1);
|
||||
|
||||
# ifdef __RAY_DIFFERENTIALS__
|
||||
if (dx)
|
||||
@@ -96,7 +96,7 @@ ccl_device float2 curve_attribute_float2(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float2, offset);
|
||||
return kernel_data_fetch(attributes_float2, offset);
|
||||
}
|
||||
else {
|
||||
return make_float2(0.0f, 0.0f);
|
||||
@@ -111,12 +111,12 @@ ccl_device float3 curve_attribute_float3(KernelGlobals kg,
|
||||
ccl_private float3 *dy)
|
||||
{
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
|
||||
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
|
||||
KernelCurve curve = kernel_data_fetch(curves, sd->prim);
|
||||
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
|
||||
float3 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
|
||||
float3 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + k1);
|
||||
float3 f0 = kernel_data_fetch(attributes_float3, desc.offset + k0);
|
||||
float3 f1 = kernel_data_fetch(attributes_float3, desc.offset + k1);
|
||||
|
||||
# ifdef __RAY_DIFFERENTIALS__
|
||||
if (dx)
|
||||
@@ -138,7 +138,7 @@ ccl_device float3 curve_attribute_float3(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float3, offset);
|
||||
return kernel_data_fetch(attributes_float3, offset);
|
||||
}
|
||||
else {
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
@@ -153,12 +153,12 @@ ccl_device float4 curve_attribute_float4(KernelGlobals kg,
|
||||
ccl_private float4 *dy)
|
||||
{
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE_KEY | ATTR_ELEMENT_CURVE_KEY_MOTION)) {
|
||||
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
|
||||
KernelCurve curve = kernel_data_fetch(curves, sd->prim);
|
||||
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
|
||||
float4 f0 = kernel_tex_fetch(__attributes_float4, desc.offset + k0);
|
||||
float4 f1 = kernel_tex_fetch(__attributes_float4, desc.offset + k1);
|
||||
float4 f0 = kernel_data_fetch(attributes_float4, desc.offset + k0);
|
||||
float4 f1 = kernel_data_fetch(attributes_float4, desc.offset + k1);
|
||||
|
||||
# ifdef __RAY_DIFFERENTIALS__
|
||||
if (dx)
|
||||
@@ -180,7 +180,7 @@ ccl_device float4 curve_attribute_float4(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float4, offset);
|
||||
return kernel_data_fetch(attributes_float4, offset);
|
||||
}
|
||||
else {
|
||||
return zero_float4();
|
||||
@@ -195,15 +195,15 @@ ccl_device float curve_thickness(KernelGlobals kg, ccl_private const ShaderData
|
||||
float r = 0.0f;
|
||||
|
||||
if (sd->type & PRIMITIVE_CURVE) {
|
||||
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
|
||||
KernelCurve curve = kernel_data_fetch(curves, sd->prim);
|
||||
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
|
||||
float4 P_curve[2];
|
||||
|
||||
if (!(sd->type & PRIMITIVE_MOTION)) {
|
||||
P_curve[0] = kernel_tex_fetch(__curve_keys, k0);
|
||||
P_curve[1] = kernel_tex_fetch(__curve_keys, k1);
|
||||
P_curve[0] = kernel_data_fetch(curve_keys, k0);
|
||||
P_curve[1] = kernel_data_fetch(curve_keys, k1);
|
||||
}
|
||||
else {
|
||||
motion_curve_keys_linear(kg, sd->object, sd->prim, sd->time, k0, k1, P_curve);
|
||||
@@ -232,14 +232,14 @@ ccl_device float curve_random(KernelGlobals kg, ccl_private const ShaderData *sd
|
||||
|
||||
ccl_device float3 curve_motion_center_location(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
KernelCurve curve = kernel_tex_fetch(__curves, sd->prim);
|
||||
KernelCurve curve = kernel_data_fetch(curves, sd->prim);
|
||||
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
|
||||
float4 P_curve[2];
|
||||
|
||||
P_curve[0] = kernel_tex_fetch(__curve_keys, k0);
|
||||
P_curve[1] = kernel_tex_fetch(__curve_keys, k1);
|
||||
P_curve[0] = kernel_data_fetch(curve_keys, k0);
|
||||
P_curve[1] = kernel_data_fetch(curve_keys, k1);
|
||||
|
||||
return float4_to_float3(P_curve[1]) * sd->u + float4_to_float3(P_curve[0]) * (1.0f - sd->u);
|
||||
}
|
||||
|
@@ -624,7 +624,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
|
||||
{
|
||||
const bool is_motion = (type & PRIMITIVE_MOTION);
|
||||
|
||||
KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
|
||||
KernelCurve kcurve = kernel_data_fetch(curves, prim);
|
||||
|
||||
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
|
||||
int k1 = k0 + 1;
|
||||
@@ -633,10 +633,10 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
|
||||
|
||||
float4 curve[4];
|
||||
if (!is_motion) {
|
||||
curve[0] = kernel_tex_fetch(__curve_keys, ka);
|
||||
curve[1] = kernel_tex_fetch(__curve_keys, k0);
|
||||
curve[2] = kernel_tex_fetch(__curve_keys, k1);
|
||||
curve[3] = kernel_tex_fetch(__curve_keys, kb);
|
||||
curve[0] = kernel_data_fetch(curve_keys, ka);
|
||||
curve[1] = kernel_data_fetch(curve_keys, k0);
|
||||
curve[2] = kernel_data_fetch(curve_keys, k1);
|
||||
curve[3] = kernel_data_fetch(curve_keys, kb);
|
||||
}
|
||||
else {
|
||||
motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
|
||||
@@ -682,7 +682,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
|
||||
D = safe_normalize_len(D, &t);
|
||||
}
|
||||
|
||||
KernelCurve kcurve = kernel_tex_fetch(__curves, isect_prim);
|
||||
KernelCurve kcurve = kernel_data_fetch(curves, isect_prim);
|
||||
|
||||
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
@@ -692,10 +692,10 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
|
||||
float4 P_curve[4];
|
||||
|
||||
if (!(sd->type & PRIMITIVE_MOTION)) {
|
||||
P_curve[0] = kernel_tex_fetch(__curve_keys, ka);
|
||||
P_curve[1] = kernel_tex_fetch(__curve_keys, k0);
|
||||
P_curve[2] = kernel_tex_fetch(__curve_keys, k1);
|
||||
P_curve[3] = kernel_tex_fetch(__curve_keys, kb);
|
||||
P_curve[0] = kernel_data_fetch(curve_keys, ka);
|
||||
P_curve[1] = kernel_data_fetch(curve_keys, k0);
|
||||
P_curve[2] = kernel_data_fetch(curve_keys, k1);
|
||||
P_curve[3] = kernel_data_fetch(curve_keys, kb);
|
||||
}
|
||||
else {
|
||||
motion_curve_keys(kg, sd->object, sd->prim, sd->time, ka, k0, k1, kb, P_curve);
|
||||
@@ -750,7 +750,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg,
|
||||
sd->P = P;
|
||||
sd->Ng = (sd->type & PRIMITIVE_CURVE_RIBBON) ? sd->I : sd->N;
|
||||
sd->dPdv = cross(sd->dPdu, sd->Ng);
|
||||
sd->shader = kernel_tex_fetch(__curves, sd->prim).shader_id;
|
||||
sd->shader = kernel_data_fetch(curves, sd->prim).shader_id;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
@@ -27,8 +27,8 @@ ccl_device_inline void motion_curve_keys_for_step_linear(KernelGlobals kg,
|
||||
{
|
||||
if (step == numsteps) {
|
||||
/* center step: regular key location */
|
||||
keys[0] = kernel_tex_fetch(__curve_keys, k0);
|
||||
keys[1] = kernel_tex_fetch(__curve_keys, k1);
|
||||
keys[0] = kernel_data_fetch(curve_keys, k0);
|
||||
keys[1] = kernel_data_fetch(curve_keys, k1);
|
||||
}
|
||||
else {
|
||||
/* center step is not stored in this array */
|
||||
@@ -37,8 +37,8 @@ ccl_device_inline void motion_curve_keys_for_step_linear(KernelGlobals kg,
|
||||
|
||||
offset += step * numkeys;
|
||||
|
||||
keys[0] = kernel_tex_fetch(__attributes_float4, offset + k0);
|
||||
keys[1] = kernel_tex_fetch(__attributes_float4, offset + k1);
|
||||
keys[0] = kernel_data_fetch(attributes_float4, offset + k0);
|
||||
keys[1] = kernel_data_fetch(attributes_float4, offset + k1);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -83,10 +83,10 @@ ccl_device_inline void motion_curve_keys_for_step(KernelGlobals kg,
|
||||
{
|
||||
if (step == numsteps) {
|
||||
/* center step: regular key location */
|
||||
keys[0] = kernel_tex_fetch(__curve_keys, k0);
|
||||
keys[1] = kernel_tex_fetch(__curve_keys, k1);
|
||||
keys[2] = kernel_tex_fetch(__curve_keys, k2);
|
||||
keys[3] = kernel_tex_fetch(__curve_keys, k3);
|
||||
keys[0] = kernel_data_fetch(curve_keys, k0);
|
||||
keys[1] = kernel_data_fetch(curve_keys, k1);
|
||||
keys[2] = kernel_data_fetch(curve_keys, k2);
|
||||
keys[3] = kernel_data_fetch(curve_keys, k3);
|
||||
}
|
||||
else {
|
||||
/* center step is not stored in this array */
|
||||
@@ -95,10 +95,10 @@ ccl_device_inline void motion_curve_keys_for_step(KernelGlobals kg,
|
||||
|
||||
offset += step * numkeys;
|
||||
|
||||
keys[0] = kernel_tex_fetch(__attributes_float4, offset + k0);
|
||||
keys[1] = kernel_tex_fetch(__attributes_float4, offset + k1);
|
||||
keys[2] = kernel_tex_fetch(__attributes_float4, offset + k2);
|
||||
keys[3] = kernel_tex_fetch(__attributes_float4, offset + k3);
|
||||
keys[0] = kernel_data_fetch(attributes_float4, offset + k0);
|
||||
keys[1] = kernel_data_fetch(attributes_float4, offset + k1);
|
||||
keys[2] = kernel_data_fetch(attributes_float4, offset + k2);
|
||||
keys[3] = kernel_data_fetch(attributes_float4, offset + k3);
|
||||
}
|
||||
}
|
||||
|
||||
|
@@ -19,7 +19,7 @@ motion_point_for_step(KernelGlobals kg, int offset, int numkeys, int numsteps, i
|
||||
{
|
||||
if (step == numsteps) {
|
||||
/* center step: regular key location */
|
||||
return kernel_tex_fetch(__points, prim);
|
||||
return kernel_data_fetch(points, prim);
|
||||
}
|
||||
else {
|
||||
/* center step is not stored in this array */
|
||||
@@ -28,7 +28,7 @@ motion_point_for_step(KernelGlobals kg, int offset, int numkeys, int numsteps, i
|
||||
|
||||
offset += step * numkeys;
|
||||
|
||||
return kernel_tex_fetch(__attributes_float4, offset + prim);
|
||||
return kernel_data_fetch(attributes_float4, offset + prim);
|
||||
}
|
||||
}
|
||||
|
||||
|
@@ -30,9 +30,9 @@ ccl_device_inline void motion_triangle_verts_for_step(KernelGlobals kg,
|
||||
{
|
||||
if (step == numsteps) {
|
||||
/* center step: regular vertex location */
|
||||
verts[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
|
||||
verts[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
|
||||
verts[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
|
||||
verts[0] = kernel_data_fetch(tri_verts, tri_vindex.w + 0);
|
||||
verts[1] = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
|
||||
verts[2] = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
|
||||
}
|
||||
else {
|
||||
/* center step not store in this array */
|
||||
@@ -41,9 +41,9 @@ ccl_device_inline void motion_triangle_verts_for_step(KernelGlobals kg,
|
||||
|
||||
offset += step * numverts;
|
||||
|
||||
verts[0] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x);
|
||||
verts[1] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y);
|
||||
verts[2] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z);
|
||||
verts[0] = kernel_data_fetch(attributes_float3, offset + tri_vindex.x);
|
||||
verts[1] = kernel_data_fetch(attributes_float3, offset + tri_vindex.y);
|
||||
verts[2] = kernel_data_fetch(attributes_float3, offset + tri_vindex.z);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -57,9 +57,9 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals kg,
|
||||
{
|
||||
if (step == numsteps) {
|
||||
/* center step: regular vertex location */
|
||||
normals[0] = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
|
||||
normals[1] = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
|
||||
normals[2] = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
|
||||
normals[0] = kernel_data_fetch(tri_vnormal, tri_vindex.x);
|
||||
normals[1] = kernel_data_fetch(tri_vnormal, tri_vindex.y);
|
||||
normals[2] = kernel_data_fetch(tri_vnormal, tri_vindex.z);
|
||||
}
|
||||
else {
|
||||
/* center step is not stored in this array */
|
||||
@@ -68,9 +68,9 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals kg,
|
||||
|
||||
offset += step * numverts;
|
||||
|
||||
normals[0] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x);
|
||||
normals[1] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y);
|
||||
normals[2] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z);
|
||||
normals[0] = kernel_data_fetch(attributes_float3, offset + tri_vindex.x);
|
||||
normals[1] = kernel_data_fetch(attributes_float3, offset + tri_vindex.y);
|
||||
normals[2] = kernel_data_fetch(attributes_float3, offset + tri_vindex.z);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -92,7 +92,7 @@ ccl_device_inline void motion_triangle_vertices(
|
||||
|
||||
/* fetch vertex coordinates */
|
||||
float3 next_verts[3];
|
||||
uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
|
||||
motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step, verts);
|
||||
motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step + 1, next_verts);
|
||||
@@ -121,7 +121,7 @@ ccl_device_inline void motion_triangle_vertices_and_normals(
|
||||
|
||||
/* Fetch vertex coordinates. */
|
||||
float3 next_verts[3];
|
||||
uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
|
||||
motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step, verts);
|
||||
motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step + 1, next_verts);
|
||||
@@ -167,7 +167,7 @@ ccl_device_inline float3 motion_triangle_smooth_normal(
|
||||
|
||||
/* fetch normals */
|
||||
float3 normals[3], next_normals[3];
|
||||
uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
|
||||
motion_triangle_normals_for_step(kg, tri_vindex, offset, numverts, numsteps, step, normals);
|
||||
motion_triangle_normals_for_step(
|
||||
|
@@ -63,7 +63,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
|
||||
/* Visibility flag test. we do it here under the assumption
|
||||
* that most triangles are culled by node flags.
|
||||
*/
|
||||
if (kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
|
||||
if (kernel_data_fetch(prim_visibility, prim_addr) & visibility)
|
||||
#endif
|
||||
{
|
||||
isect->t = t;
|
||||
|
@@ -31,7 +31,7 @@ ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals kg,
|
||||
bool is_local)
|
||||
{
|
||||
/* Get shader. */
|
||||
sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
|
||||
sd->shader = kernel_data_fetch(tri_shader, sd->prim);
|
||||
/* Get motion info. */
|
||||
/* TODO(sergey): This logic is really similar to motion_triangle_vertices(),
|
||||
* can we de-duplicate something here?
|
||||
@@ -47,7 +47,7 @@ ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals kg,
|
||||
kernel_assert(offset != ATTR_STD_NOT_FOUND);
|
||||
/* Fetch vertex coordinates. */
|
||||
float3 verts[3], next_verts[3];
|
||||
uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step, verts);
|
||||
motion_triangle_verts_for_step(kg, tri_vindex, offset, numverts, numsteps, step + 1, next_verts);
|
||||
/* Interpolate between steps. */
|
||||
|
@@ -31,10 +31,10 @@ ccl_device_inline Transform object_fetch_transform(KernelGlobals kg,
|
||||
enum ObjectTransform type)
|
||||
{
|
||||
if (type == OBJECT_INVERSE_TRANSFORM) {
|
||||
return kernel_tex_fetch(__objects, object).itfm;
|
||||
return kernel_data_fetch(objects, object).itfm;
|
||||
}
|
||||
else {
|
||||
return kernel_tex_fetch(__objects, object).tfm;
|
||||
return kernel_data_fetch(objects, object).tfm;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -43,10 +43,10 @@ ccl_device_inline Transform object_fetch_transform(KernelGlobals kg,
|
||||
ccl_device_inline Transform lamp_fetch_transform(KernelGlobals kg, int lamp, bool inverse)
|
||||
{
|
||||
if (inverse) {
|
||||
return kernel_tex_fetch(__lights, lamp).itfm;
|
||||
return kernel_data_fetch(lights, lamp).itfm;
|
||||
}
|
||||
else {
|
||||
return kernel_tex_fetch(__lights, lamp).tfm;
|
||||
return kernel_data_fetch(lights, lamp).tfm;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -57,7 +57,7 @@ ccl_device_inline Transform object_fetch_motion_pass_transform(KernelGlobals kg,
|
||||
enum ObjectVectorTransform type)
|
||||
{
|
||||
int offset = object * OBJECT_MOTION_PASS_SIZE + (int)type;
|
||||
return kernel_tex_fetch(__object_motion_pass, offset);
|
||||
return kernel_data_fetch(object_motion_pass, offset);
|
||||
}
|
||||
|
||||
/* Motion blurred object transformations */
|
||||
@@ -65,9 +65,9 @@ ccl_device_inline Transform object_fetch_motion_pass_transform(KernelGlobals kg,
|
||||
#ifdef __OBJECT_MOTION__
|
||||
ccl_device_inline Transform object_fetch_transform_motion(KernelGlobals kg, int object, float time)
|
||||
{
|
||||
const uint motion_offset = kernel_tex_fetch(__objects, object).motion_offset;
|
||||
ccl_global const DecomposedTransform *motion = &kernel_tex_fetch(__object_motion, motion_offset);
|
||||
const uint num_steps = kernel_tex_fetch(__objects, object).numsteps * 2 + 1;
|
||||
const uint motion_offset = kernel_data_fetch(objects, object).motion_offset;
|
||||
ccl_global const DecomposedTransform *motion = &kernel_data_fetch(object_motion, motion_offset);
|
||||
const uint num_steps = kernel_data_fetch(objects, object).numsteps * 2 + 1;
|
||||
|
||||
Transform tfm;
|
||||
transform_motion_array_interpolate(&tfm, motion, num_steps, time);
|
||||
@@ -80,7 +80,7 @@ ccl_device_inline Transform object_fetch_transform_motion_test(KernelGlobals kg,
|
||||
float time,
|
||||
ccl_private Transform *itfm)
|
||||
{
|
||||
int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (object_flag & SD_OBJECT_MOTION) {
|
||||
/* if we do motion blur */
|
||||
Transform tfm = object_fetch_transform_motion(kg, object, time);
|
||||
@@ -259,7 +259,7 @@ ccl_device_inline float3 object_color(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
|
||||
ccl_global const KernelObject *kobject = &kernel_tex_fetch(__objects, object);
|
||||
ccl_global const KernelObject *kobject = &kernel_data_fetch(objects, object);
|
||||
return make_float3(kobject->color[0], kobject->color[1], kobject->color[2]);
|
||||
}
|
||||
|
||||
@@ -270,7 +270,7 @@ ccl_device_inline float object_alpha(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).alpha;
|
||||
return kernel_data_fetch(objects, object).alpha;
|
||||
}
|
||||
|
||||
/* Pass ID number of object */
|
||||
@@ -280,7 +280,7 @@ ccl_device_inline float object_pass_id(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).pass_id;
|
||||
return kernel_data_fetch(objects, object).pass_id;
|
||||
}
|
||||
|
||||
/* Lightgroup of lamp */
|
||||
@@ -290,7 +290,7 @@ ccl_device_inline int lamp_lightgroup(KernelGlobals kg, int lamp)
|
||||
if (lamp == LAMP_NONE)
|
||||
return LIGHTGROUP_NONE;
|
||||
|
||||
return kernel_tex_fetch(__lights, lamp).lightgroup;
|
||||
return kernel_data_fetch(lights, lamp).lightgroup;
|
||||
}
|
||||
|
||||
/* Lightgroup of object */
|
||||
@@ -300,7 +300,7 @@ ccl_device_inline int object_lightgroup(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return LIGHTGROUP_NONE;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).lightgroup;
|
||||
return kernel_data_fetch(objects, object).lightgroup;
|
||||
}
|
||||
|
||||
/* Per lamp random number for shader variation */
|
||||
@@ -310,7 +310,7 @@ ccl_device_inline float lamp_random_number(KernelGlobals kg, int lamp)
|
||||
if (lamp == LAMP_NONE)
|
||||
return 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__lights, lamp).random;
|
||||
return kernel_data_fetch(lights, lamp).random;
|
||||
}
|
||||
|
||||
/* Per object random number for shader variation */
|
||||
@@ -320,7 +320,7 @@ ccl_device_inline float object_random_number(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).random_number;
|
||||
return kernel_data_fetch(objects, object).random_number;
|
||||
}
|
||||
|
||||
/* Particle ID from which this object was generated */
|
||||
@@ -330,7 +330,7 @@ ccl_device_inline int object_particle_id(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return 0;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).particle_index;
|
||||
return kernel_data_fetch(objects, object).particle_index;
|
||||
}
|
||||
|
||||
/* Generated texture coordinate on surface from where object was instanced */
|
||||
@@ -340,7 +340,7 @@ ccl_device_inline float3 object_dupli_generated(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
|
||||
ccl_global const KernelObject *kobject = &kernel_tex_fetch(__objects, object);
|
||||
ccl_global const KernelObject *kobject = &kernel_data_fetch(objects, object);
|
||||
return make_float3(
|
||||
kobject->dupli_generated[0], kobject->dupli_generated[1], kobject->dupli_generated[2]);
|
||||
}
|
||||
@@ -352,7 +352,7 @@ ccl_device_inline float3 object_dupli_uv(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
|
||||
ccl_global const KernelObject *kobject = &kernel_tex_fetch(__objects, object);
|
||||
ccl_global const KernelObject *kobject = &kernel_data_fetch(objects, object);
|
||||
return make_float3(kobject->dupli_uv[0], kobject->dupli_uv[1], 0.0f);
|
||||
}
|
||||
|
||||
@@ -365,13 +365,13 @@ ccl_device_inline void object_motion_info(KernelGlobals kg,
|
||||
ccl_private int *numkeys)
|
||||
{
|
||||
if (numkeys) {
|
||||
*numkeys = kernel_tex_fetch(__objects, object).numkeys;
|
||||
*numkeys = kernel_data_fetch(objects, object).numkeys;
|
||||
}
|
||||
|
||||
if (numsteps)
|
||||
*numsteps = kernel_tex_fetch(__objects, object).numsteps;
|
||||
*numsteps = kernel_data_fetch(objects, object).numsteps;
|
||||
if (numverts)
|
||||
*numverts = kernel_tex_fetch(__objects, object).numverts;
|
||||
*numverts = kernel_data_fetch(objects, object).numverts;
|
||||
}
|
||||
|
||||
/* Offset to an objects patch map */
|
||||
@@ -381,7 +381,7 @@ ccl_device_inline uint object_patch_map_offset(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return 0;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).patch_map_offset;
|
||||
return kernel_data_fetch(objects, object).patch_map_offset;
|
||||
}
|
||||
|
||||
/* Volume step size */
|
||||
@@ -392,7 +392,7 @@ ccl_device_inline float object_volume_density(KernelGlobals kg, int object)
|
||||
return 1.0f;
|
||||
}
|
||||
|
||||
return kernel_tex_fetch(__objects, object).volume_density;
|
||||
return kernel_data_fetch(objects, object).volume_density;
|
||||
}
|
||||
|
||||
ccl_device_inline float object_volume_step_size(KernelGlobals kg, int object)
|
||||
@@ -401,14 +401,14 @@ ccl_device_inline float object_volume_step_size(KernelGlobals kg, int object)
|
||||
return kernel_data.background.volume_step_size;
|
||||
}
|
||||
|
||||
return kernel_tex_fetch(__object_volume_step, object);
|
||||
return kernel_data_fetch(object_volume_step, object);
|
||||
}
|
||||
|
||||
/* Pass ID for shader */
|
||||
|
||||
ccl_device int shader_pass_id(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
return kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).pass_id;
|
||||
return kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).pass_id;
|
||||
}
|
||||
|
||||
/* Cryptomatte ID */
|
||||
@@ -418,7 +418,7 @@ ccl_device_inline float object_cryptomatte_id(KernelGlobals kg, int object)
|
||||
if (object == OBJECT_NONE)
|
||||
return 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).cryptomatte_object;
|
||||
return kernel_data_fetch(objects, object).cryptomatte_object;
|
||||
}
|
||||
|
||||
ccl_device_inline float object_cryptomatte_asset_id(KernelGlobals kg, int object)
|
||||
@@ -426,49 +426,49 @@ ccl_device_inline float object_cryptomatte_asset_id(KernelGlobals kg, int object
|
||||
if (object == OBJECT_NONE)
|
||||
return 0;
|
||||
|
||||
return kernel_tex_fetch(__objects, object).cryptomatte_asset;
|
||||
return kernel_data_fetch(objects, object).cryptomatte_asset;
|
||||
}
|
||||
|
||||
/* Particle data from which object was instanced */
|
||||
|
||||
ccl_device_inline uint particle_index(KernelGlobals kg, int particle)
|
||||
{
|
||||
return kernel_tex_fetch(__particles, particle).index;
|
||||
return kernel_data_fetch(particles, particle).index;
|
||||
}
|
||||
|
||||
ccl_device float particle_age(KernelGlobals kg, int particle)
|
||||
{
|
||||
return kernel_tex_fetch(__particles, particle).age;
|
||||
return kernel_data_fetch(particles, particle).age;
|
||||
}
|
||||
|
||||
ccl_device float particle_lifetime(KernelGlobals kg, int particle)
|
||||
{
|
||||
return kernel_tex_fetch(__particles, particle).lifetime;
|
||||
return kernel_data_fetch(particles, particle).lifetime;
|
||||
}
|
||||
|
||||
ccl_device float particle_size(KernelGlobals kg, int particle)
|
||||
{
|
||||
return kernel_tex_fetch(__particles, particle).size;
|
||||
return kernel_data_fetch(particles, particle).size;
|
||||
}
|
||||
|
||||
ccl_device float4 particle_rotation(KernelGlobals kg, int particle)
|
||||
{
|
||||
return kernel_tex_fetch(__particles, particle).rotation;
|
||||
return kernel_data_fetch(particles, particle).rotation;
|
||||
}
|
||||
|
||||
ccl_device float3 particle_location(KernelGlobals kg, int particle)
|
||||
{
|
||||
return float4_to_float3(kernel_tex_fetch(__particles, particle).location);
|
||||
return float4_to_float3(kernel_data_fetch(particles, particle).location);
|
||||
}
|
||||
|
||||
ccl_device float3 particle_velocity(KernelGlobals kg, int particle)
|
||||
{
|
||||
return float4_to_float3(kernel_tex_fetch(__particles, particle).velocity);
|
||||
return float4_to_float3(kernel_data_fetch(particles, particle).velocity);
|
||||
}
|
||||
|
||||
ccl_device float3 particle_angular_velocity(KernelGlobals kg, int particle)
|
||||
{
|
||||
return float4_to_float3(kernel_tex_fetch(__particles, particle).angular_velocity);
|
||||
return float4_to_float3(kernel_data_fetch(particles, particle).angular_velocity);
|
||||
}
|
||||
|
||||
/* Object intersection in BVH */
|
||||
|
@@ -62,7 +62,7 @@ patch_map_find_patch(KernelGlobals kg, int object, int patch, float u, float v)
|
||||
int quadrant = patch_map_resolve_quadrant(median, &u, &v);
|
||||
kernel_assert(quadrant >= 0);
|
||||
|
||||
uint child = kernel_tex_fetch(__patches, node + quadrant);
|
||||
uint child = kernel_data_fetch(patches, node + quadrant);
|
||||
|
||||
/* is the quadrant a hole? */
|
||||
if (!(child & PATCH_MAP_NODE_IS_SET)) {
|
||||
@@ -73,9 +73,9 @@ patch_map_find_patch(KernelGlobals kg, int object, int patch, float u, float v)
|
||||
uint index = child & PATCH_MAP_NODE_INDEX_MASK;
|
||||
|
||||
if (child & PATCH_MAP_NODE_IS_LEAF) {
|
||||
handle.array_index = kernel_tex_fetch(__patches, index + 0);
|
||||
handle.patch_index = kernel_tex_fetch(__patches, index + 1);
|
||||
handle.vert_index = kernel_tex_fetch(__patches, index + 2);
|
||||
handle.array_index = kernel_data_fetch(patches, index + 0);
|
||||
handle.patch_index = kernel_data_fetch(patches, index + 1);
|
||||
handle.vert_index = kernel_data_fetch(patches, index + 2);
|
||||
|
||||
return handle;
|
||||
}
|
||||
@@ -189,11 +189,11 @@ ccl_device_inline int patch_eval_indices(KernelGlobals kg,
|
||||
int channel,
|
||||
int indices[PATCH_MAX_CONTROL_VERTS])
|
||||
{
|
||||
int index_base = kernel_tex_fetch(__patches, handle->array_index + 2) + handle->vert_index;
|
||||
int index_base = kernel_data_fetch(patches, handle->array_index + 2) + handle->vert_index;
|
||||
|
||||
/* XXX: regular patches only */
|
||||
for (int i = 0; i < 16; i++) {
|
||||
indices[i] = kernel_tex_fetch(__patches, index_base + i);
|
||||
indices[i] = kernel_data_fetch(patches, index_base + i);
|
||||
}
|
||||
|
||||
return 16;
|
||||
@@ -209,7 +209,7 @@ ccl_device_inline void patch_eval_basis(KernelGlobals kg,
|
||||
float weights_du[PATCH_MAX_CONTROL_VERTS],
|
||||
float weights_dv[PATCH_MAX_CONTROL_VERTS])
|
||||
{
|
||||
uint patch_bits = kernel_tex_fetch(__patches, handle->patch_index + 1); /* read patch param */
|
||||
uint patch_bits = kernel_data_fetch(patches, handle->patch_index + 1); /* read patch param */
|
||||
float d_scale = 1 << patch_eval_depth(patch_bits);
|
||||
|
||||
bool non_quad_root = (patch_bits >> 4) & 0x1;
|
||||
@@ -287,7 +287,7 @@ ccl_device float patch_eval_float(KernelGlobals kg,
|
||||
*dv = 0.0f;
|
||||
|
||||
for (int i = 0; i < num_control; i++) {
|
||||
float v = kernel_tex_fetch(__attributes_float, offset + indices[i]);
|
||||
float v = kernel_data_fetch(attributes_float, offset + indices[i]);
|
||||
|
||||
val += v * weights[i];
|
||||
if (du)
|
||||
@@ -324,7 +324,7 @@ ccl_device float2 patch_eval_float2(KernelGlobals kg,
|
||||
*dv = make_float2(0.0f, 0.0f);
|
||||
|
||||
for (int i = 0; i < num_control; i++) {
|
||||
float2 v = kernel_tex_fetch(__attributes_float2, offset + indices[i]);
|
||||
float2 v = kernel_data_fetch(attributes_float2, offset + indices[i]);
|
||||
|
||||
val += v * weights[i];
|
||||
if (du)
|
||||
@@ -361,7 +361,7 @@ ccl_device float3 patch_eval_float3(KernelGlobals kg,
|
||||
*dv = make_float3(0.0f, 0.0f, 0.0f);
|
||||
|
||||
for (int i = 0; i < num_control; i++) {
|
||||
float3 v = kernel_tex_fetch(__attributes_float3, offset + indices[i]);
|
||||
float3 v = kernel_data_fetch(attributes_float3, offset + indices[i]);
|
||||
|
||||
val += v * weights[i];
|
||||
if (du)
|
||||
@@ -398,7 +398,7 @@ ccl_device float4 patch_eval_float4(KernelGlobals kg,
|
||||
*dv = zero_float4();
|
||||
|
||||
for (int i = 0; i < num_control; i++) {
|
||||
float4 v = kernel_tex_fetch(__attributes_float4, offset + indices[i]);
|
||||
float4 v = kernel_data_fetch(attributes_float4, offset + indices[i]);
|
||||
|
||||
val += v * weights[i];
|
||||
if (du)
|
||||
@@ -436,7 +436,7 @@ ccl_device float4 patch_eval_uchar4(KernelGlobals kg,
|
||||
|
||||
for (int i = 0; i < num_control; i++) {
|
||||
float4 v = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, offset + indices[i])));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, offset + indices[i])));
|
||||
|
||||
val += v * weights[i];
|
||||
if (du)
|
||||
|
@@ -26,7 +26,7 @@ ccl_device float point_attribute_float(KernelGlobals kg,
|
||||
# endif
|
||||
|
||||
if (desc.element == ATTR_ELEMENT_VERTEX) {
|
||||
return kernel_tex_fetch(__attributes_float, desc.offset + sd->prim);
|
||||
return kernel_data_fetch(attributes_float, desc.offset + sd->prim);
|
||||
}
|
||||
else {
|
||||
return 0.0f;
|
||||
@@ -47,7 +47,7 @@ ccl_device float2 point_attribute_float2(KernelGlobals kg,
|
||||
# endif
|
||||
|
||||
if (desc.element == ATTR_ELEMENT_VERTEX) {
|
||||
return kernel_tex_fetch(__attributes_float2, desc.offset + sd->prim);
|
||||
return kernel_data_fetch(attributes_float2, desc.offset + sd->prim);
|
||||
}
|
||||
else {
|
||||
return make_float2(0.0f, 0.0f);
|
||||
@@ -68,7 +68,7 @@ ccl_device float3 point_attribute_float3(KernelGlobals kg,
|
||||
# endif
|
||||
|
||||
if (desc.element == ATTR_ELEMENT_VERTEX) {
|
||||
return kernel_tex_fetch(__attributes_float3, desc.offset + sd->prim);
|
||||
return kernel_data_fetch(attributes_float3, desc.offset + sd->prim);
|
||||
}
|
||||
else {
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
@@ -89,7 +89,7 @@ ccl_device float4 point_attribute_float4(KernelGlobals kg,
|
||||
# endif
|
||||
|
||||
if (desc.element == ATTR_ELEMENT_VERTEX) {
|
||||
return kernel_tex_fetch(__attributes_float4, desc.offset + sd->prim);
|
||||
return kernel_data_fetch(attributes_float4, desc.offset + sd->prim);
|
||||
}
|
||||
else {
|
||||
return zero_float4();
|
||||
@@ -104,7 +104,7 @@ ccl_device float3 point_position(KernelGlobals kg, ccl_private const ShaderData
|
||||
/* World space center. */
|
||||
float3 P = (sd->type & PRIMITIVE_MOTION) ?
|
||||
float4_to_float3(motion_point(kg, sd->object, sd->prim, sd->time)) :
|
||||
float4_to_float3(kernel_tex_fetch(__points, sd->prim));
|
||||
float4_to_float3(kernel_data_fetch(points, sd->prim));
|
||||
|
||||
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
|
||||
object_position_transform(kg, sd, &P);
|
||||
@@ -122,7 +122,7 @@ ccl_device float point_radius(KernelGlobals kg, ccl_private const ShaderData *sd
|
||||
{
|
||||
if (sd->type & PRIMITIVE_POINT) {
|
||||
/* World space radius. */
|
||||
const float r = kernel_tex_fetch(__points, sd->prim).w;
|
||||
const float r = kernel_data_fetch(points, sd->prim).w;
|
||||
|
||||
if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
|
||||
return r;
|
||||
@@ -155,7 +155,7 @@ ccl_device float point_random(KernelGlobals kg, ccl_private const ShaderData *sd
|
||||
|
||||
ccl_device float3 point_motion_center_location(KernelGlobals kg, ccl_private const ShaderData *sd)
|
||||
{
|
||||
return float4_to_float3(kernel_tex_fetch(__points, sd->prim));
|
||||
return float4_to_float3(kernel_data_fetch(points, sd->prim));
|
||||
}
|
||||
|
||||
#endif /* __POINTCLOUD__ */
|
||||
|
@@ -63,7 +63,7 @@ ccl_device_forceinline bool point_intersect(KernelGlobals kg,
|
||||
const int type)
|
||||
{
|
||||
const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) :
|
||||
kernel_tex_fetch(__points, prim);
|
||||
kernel_data_fetch(points, prim);
|
||||
|
||||
if (!point_intersect_test(point, P, dir, tmax, &isect->t)) {
|
||||
return false;
|
||||
@@ -82,7 +82,7 @@ ccl_device_inline void point_shader_setup(KernelGlobals kg,
|
||||
ccl_private const Intersection *isect,
|
||||
ccl_private const Ray *ray)
|
||||
{
|
||||
sd->shader = kernel_tex_fetch(__points_shader, isect->prim);
|
||||
sd->shader = kernel_data_fetch(points_shader, isect->prim);
|
||||
sd->P = ray->P + ray->D * isect->t;
|
||||
|
||||
/* Texture coordinates, zero for now. */
|
||||
@@ -94,7 +94,7 @@ ccl_device_inline void point_shader_setup(KernelGlobals kg,
|
||||
/* Compute point center for normal. */
|
||||
float3 center = float4_to_float3((isect->type & PRIMITIVE_MOTION) ?
|
||||
motion_point(kg, sd->object, sd->prim, sd->time) :
|
||||
kernel_tex_fetch(__points, sd->prim));
|
||||
kernel_data_fetch(points, sd->prim));
|
||||
if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
|
||||
object_position_transform_auto(kg, sd, ¢er);
|
||||
}
|
||||
|
@@ -40,7 +40,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
|
||||
sd->ray_length = isect->t;
|
||||
sd->type = isect->type;
|
||||
sd->object = isect->object;
|
||||
sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
|
||||
sd->object_flag = kernel_data_fetch(object_flag, sd->object);
|
||||
sd->prim = isect->prim;
|
||||
sd->lamp = LAMP_NONE;
|
||||
sd->flag = 0;
|
||||
@@ -73,7 +73,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
|
||||
if (sd->type == PRIMITIVE_TRIANGLE) {
|
||||
/* static triangle */
|
||||
float3 Ng = triangle_normal(kg, sd);
|
||||
sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
|
||||
sd->shader = kernel_data_fetch(tri_shader, sd->prim);
|
||||
|
||||
/* vectors */
|
||||
sd->P = triangle_point_from_uv(kg, sd, isect->object, isect->prim, isect->u, isect->v);
|
||||
@@ -106,7 +106,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg,
|
||||
}
|
||||
}
|
||||
|
||||
sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->flag = kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
|
||||
/* backfacing test */
|
||||
bool backfacing = (dot(sd->Ng, sd->I) < 0.0f);
|
||||
@@ -169,10 +169,10 @@ ccl_device_inline void shader_setup_from_sample(KernelGlobals kg,
|
||||
sd->time = time;
|
||||
sd->ray_length = t;
|
||||
|
||||
sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->flag = kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->object_flag = 0;
|
||||
if (sd->object != OBJECT_NONE) {
|
||||
sd->object_flag |= kernel_tex_fetch(__object_flag, sd->object);
|
||||
sd->object_flag |= kernel_data_fetch(object_flag, sd->object);
|
||||
|
||||
#ifdef __OBJECT_MOTION__
|
||||
shader_setup_object_transforms(kg, sd, time);
|
||||
@@ -264,21 +264,20 @@ ccl_device void shader_setup_from_displace(KernelGlobals kg,
|
||||
/* force smooth shading for displacement */
|
||||
shader |= SHADER_SMOOTH_NORMAL;
|
||||
|
||||
shader_setup_from_sample(
|
||||
kg,
|
||||
sd,
|
||||
P,
|
||||
Ng,
|
||||
I,
|
||||
shader,
|
||||
object,
|
||||
prim,
|
||||
u,
|
||||
v,
|
||||
0.0f,
|
||||
0.5f,
|
||||
!(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED),
|
||||
LAMP_NONE);
|
||||
shader_setup_from_sample(kg,
|
||||
sd,
|
||||
P,
|
||||
Ng,
|
||||
I,
|
||||
shader,
|
||||
object,
|
||||
prim,
|
||||
u,
|
||||
v,
|
||||
0.0f,
|
||||
0.5f,
|
||||
!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED),
|
||||
LAMP_NONE);
|
||||
}
|
||||
|
||||
/* ShaderData setup for point on curve. */
|
||||
@@ -300,18 +299,18 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg,
|
||||
sd->ray_length = 0.0f;
|
||||
|
||||
/* Shader */
|
||||
sd->shader = kernel_tex_fetch(__curves, prim).shader_id;
|
||||
sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->shader = kernel_data_fetch(curves, prim).shader_id;
|
||||
sd->flag = kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
|
||||
/* Object */
|
||||
sd->object = object;
|
||||
sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
|
||||
sd->object_flag = kernel_data_fetch(object_flag, sd->object);
|
||||
#ifdef __OBJECT_MOTION__
|
||||
shader_setup_object_transforms(kg, sd, sd->time);
|
||||
#endif
|
||||
|
||||
/* Get control points. */
|
||||
KernelCurve kcurve = kernel_tex_fetch(__curves, prim);
|
||||
KernelCurve kcurve = kernel_data_fetch(curves, prim);
|
||||
|
||||
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
|
||||
int k1 = k0 + 1;
|
||||
@@ -320,10 +319,10 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg,
|
||||
|
||||
float4 P_curve[4];
|
||||
|
||||
P_curve[0] = kernel_tex_fetch(__curve_keys, ka);
|
||||
P_curve[1] = kernel_tex_fetch(__curve_keys, k0);
|
||||
P_curve[2] = kernel_tex_fetch(__curve_keys, k1);
|
||||
P_curve[3] = kernel_tex_fetch(__curve_keys, kb);
|
||||
P_curve[0] = kernel_data_fetch(curve_keys, ka);
|
||||
P_curve[1] = kernel_data_fetch(curve_keys, k0);
|
||||
P_curve[2] = kernel_data_fetch(curve_keys, k1);
|
||||
P_curve[3] = kernel_data_fetch(curve_keys, kb);
|
||||
|
||||
/* Interpolate position and tangent. */
|
||||
sd->P = float4_to_float3(catmull_rom_basis_derivative(P_curve, sd->u));
|
||||
@@ -373,7 +372,7 @@ ccl_device_inline void shader_setup_from_background(KernelGlobals kg,
|
||||
sd->Ng = -ray_D;
|
||||
sd->I = -ray_D;
|
||||
sd->shader = kernel_data.background.surface_shader;
|
||||
sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->flag = kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->object_flag = 0;
|
||||
sd->time = ray_time;
|
||||
sd->ray_length = 0.0f;
|
||||
|
@@ -13,11 +13,11 @@ ccl_device_inline void subd_triangle_patch_uv(KernelGlobals kg,
|
||||
ccl_private const ShaderData *sd,
|
||||
float2 uv[3])
|
||||
{
|
||||
uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
|
||||
uv[0] = kernel_tex_fetch(__tri_patch_uv, tri_vindex.x);
|
||||
uv[1] = kernel_tex_fetch(__tri_patch_uv, tri_vindex.y);
|
||||
uv[2] = kernel_tex_fetch(__tri_patch_uv, tri_vindex.z);
|
||||
uv[0] = kernel_data_fetch(tri_patch_uv, tri_vindex.x);
|
||||
uv[1] = kernel_data_fetch(tri_patch_uv, tri_vindex.y);
|
||||
uv[2] = kernel_data_fetch(tri_patch_uv, tri_vindex.z);
|
||||
}
|
||||
|
||||
/* Vertex indices of patch */
|
||||
@@ -26,10 +26,10 @@ ccl_device_inline uint4 subd_triangle_patch_indices(KernelGlobals kg, int patch)
|
||||
{
|
||||
uint4 indices;
|
||||
|
||||
indices.x = kernel_tex_fetch(__patches, patch + 0);
|
||||
indices.y = kernel_tex_fetch(__patches, patch + 1);
|
||||
indices.z = kernel_tex_fetch(__patches, patch + 2);
|
||||
indices.w = kernel_tex_fetch(__patches, patch + 3);
|
||||
indices.x = kernel_data_fetch(patches, patch + 0);
|
||||
indices.y = kernel_data_fetch(patches, patch + 1);
|
||||
indices.z = kernel_data_fetch(patches, patch + 2);
|
||||
indices.w = kernel_data_fetch(patches, patch + 3);
|
||||
|
||||
return indices;
|
||||
}
|
||||
@@ -38,14 +38,14 @@ ccl_device_inline uint4 subd_triangle_patch_indices(KernelGlobals kg, int patch)
|
||||
|
||||
ccl_device_inline uint subd_triangle_patch_face(KernelGlobals kg, int patch)
|
||||
{
|
||||
return kernel_tex_fetch(__patches, patch + 4);
|
||||
return kernel_data_fetch(patches, patch + 4);
|
||||
}
|
||||
|
||||
/* Number of corners on originating face */
|
||||
|
||||
ccl_device_inline uint subd_triangle_patch_num_corners(KernelGlobals kg, int patch)
|
||||
{
|
||||
return kernel_tex_fetch(__patches, patch + 5) & 0xffff;
|
||||
return kernel_data_fetch(patches, patch + 5) & 0xffff;
|
||||
}
|
||||
|
||||
/* Indices of the four corners that are used by the patch */
|
||||
@@ -54,10 +54,10 @@ ccl_device_inline void subd_triangle_patch_corners(KernelGlobals kg, int patch,
|
||||
{
|
||||
uint4 data;
|
||||
|
||||
data.x = kernel_tex_fetch(__patches, patch + 4);
|
||||
data.y = kernel_tex_fetch(__patches, patch + 5);
|
||||
data.z = kernel_tex_fetch(__patches, patch + 6);
|
||||
data.w = kernel_tex_fetch(__patches, patch + 7);
|
||||
data.x = kernel_data_fetch(patches, patch + 4);
|
||||
data.y = kernel_data_fetch(patches, patch + 5);
|
||||
data.z = kernel_data_fetch(patches, patch + 6);
|
||||
data.w = kernel_data_fetch(patches, patch + 7);
|
||||
|
||||
int num_corners = data.y & 0xffff;
|
||||
|
||||
@@ -141,7 +141,7 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__attributes_float, desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
return kernel_data_fetch(attributes_float, desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
}
|
||||
else if (desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
|
||||
float2 uv[3];
|
||||
@@ -149,10 +149,10 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
|
||||
|
||||
uint4 v = subd_triangle_patch_indices(kg, patch);
|
||||
|
||||
float f0 = kernel_tex_fetch(__attributes_float, desc.offset + v.x);
|
||||
float f1 = kernel_tex_fetch(__attributes_float, desc.offset + v.y);
|
||||
float f2 = kernel_tex_fetch(__attributes_float, desc.offset + v.z);
|
||||
float f3 = kernel_tex_fetch(__attributes_float, desc.offset + v.w);
|
||||
float f0 = kernel_data_fetch(attributes_float, desc.offset + v.x);
|
||||
float f1 = kernel_data_fetch(attributes_float, desc.offset + v.y);
|
||||
float f2 = kernel_data_fetch(attributes_float, desc.offset + v.z);
|
||||
float f3 = kernel_data_fetch(attributes_float, desc.offset + v.w);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -179,10 +179,10 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
|
||||
int corners[4];
|
||||
subd_triangle_patch_corners(kg, patch, corners);
|
||||
|
||||
float f0 = kernel_tex_fetch(__attributes_float, corners[0] + desc.offset);
|
||||
float f1 = kernel_tex_fetch(__attributes_float, corners[1] + desc.offset);
|
||||
float f2 = kernel_tex_fetch(__attributes_float, corners[2] + desc.offset);
|
||||
float f3 = kernel_tex_fetch(__attributes_float, corners[3] + desc.offset);
|
||||
float f0 = kernel_data_fetch(attributes_float, corners[0] + desc.offset);
|
||||
float f1 = kernel_data_fetch(attributes_float, corners[1] + desc.offset);
|
||||
float f2 = kernel_data_fetch(attributes_float, corners[2] + desc.offset);
|
||||
float f3 = kernel_data_fetch(attributes_float, corners[3] + desc.offset);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -208,7 +208,7 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = 0.0f;
|
||||
|
||||
return kernel_tex_fetch(__attributes_float, desc.offset);
|
||||
return kernel_data_fetch(attributes_float, desc.offset);
|
||||
}
|
||||
else {
|
||||
if (dx)
|
||||
@@ -281,8 +281,7 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = make_float2(0.0f, 0.0f);
|
||||
|
||||
return kernel_tex_fetch(__attributes_float2,
|
||||
desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
return kernel_data_fetch(attributes_float2, desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
}
|
||||
else if (desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
|
||||
float2 uv[3];
|
||||
@@ -290,10 +289,10 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
|
||||
|
||||
uint4 v = subd_triangle_patch_indices(kg, patch);
|
||||
|
||||
float2 f0 = kernel_tex_fetch(__attributes_float2, desc.offset + v.x);
|
||||
float2 f1 = kernel_tex_fetch(__attributes_float2, desc.offset + v.y);
|
||||
float2 f2 = kernel_tex_fetch(__attributes_float2, desc.offset + v.z);
|
||||
float2 f3 = kernel_tex_fetch(__attributes_float2, desc.offset + v.w);
|
||||
float2 f0 = kernel_data_fetch(attributes_float2, desc.offset + v.x);
|
||||
float2 f1 = kernel_data_fetch(attributes_float2, desc.offset + v.y);
|
||||
float2 f2 = kernel_data_fetch(attributes_float2, desc.offset + v.z);
|
||||
float2 f3 = kernel_data_fetch(attributes_float2, desc.offset + v.w);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -322,10 +321,10 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
|
||||
|
||||
float2 f0, f1, f2, f3;
|
||||
|
||||
f0 = kernel_tex_fetch(__attributes_float2, corners[0] + desc.offset);
|
||||
f1 = kernel_tex_fetch(__attributes_float2, corners[1] + desc.offset);
|
||||
f2 = kernel_tex_fetch(__attributes_float2, corners[2] + desc.offset);
|
||||
f3 = kernel_tex_fetch(__attributes_float2, corners[3] + desc.offset);
|
||||
f0 = kernel_data_fetch(attributes_float2, corners[0] + desc.offset);
|
||||
f1 = kernel_data_fetch(attributes_float2, corners[1] + desc.offset);
|
||||
f2 = kernel_data_fetch(attributes_float2, corners[2] + desc.offset);
|
||||
f3 = kernel_data_fetch(attributes_float2, corners[3] + desc.offset);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -351,7 +350,7 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = make_float2(0.0f, 0.0f);
|
||||
|
||||
return kernel_tex_fetch(__attributes_float2, desc.offset);
|
||||
return kernel_data_fetch(attributes_float2, desc.offset);
|
||||
}
|
||||
else {
|
||||
if (dx)
|
||||
@@ -423,8 +422,7 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = make_float3(0.0f, 0.0f, 0.0f);
|
||||
|
||||
return kernel_tex_fetch(__attributes_float3,
|
||||
desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
return kernel_data_fetch(attributes_float3, desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
}
|
||||
else if (desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
|
||||
float2 uv[3];
|
||||
@@ -432,10 +430,10 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
|
||||
|
||||
uint4 v = subd_triangle_patch_indices(kg, patch);
|
||||
|
||||
float3 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + v.x);
|
||||
float3 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + v.y);
|
||||
float3 f2 = kernel_tex_fetch(__attributes_float3, desc.offset + v.z);
|
||||
float3 f3 = kernel_tex_fetch(__attributes_float3, desc.offset + v.w);
|
||||
float3 f0 = kernel_data_fetch(attributes_float3, desc.offset + v.x);
|
||||
float3 f1 = kernel_data_fetch(attributes_float3, desc.offset + v.y);
|
||||
float3 f2 = kernel_data_fetch(attributes_float3, desc.offset + v.z);
|
||||
float3 f3 = kernel_data_fetch(attributes_float3, desc.offset + v.w);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -464,10 +462,10 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
|
||||
|
||||
float3 f0, f1, f2, f3;
|
||||
|
||||
f0 = kernel_tex_fetch(__attributes_float3, corners[0] + desc.offset);
|
||||
f1 = kernel_tex_fetch(__attributes_float3, corners[1] + desc.offset);
|
||||
f2 = kernel_tex_fetch(__attributes_float3, corners[2] + desc.offset);
|
||||
f3 = kernel_tex_fetch(__attributes_float3, corners[3] + desc.offset);
|
||||
f0 = kernel_data_fetch(attributes_float3, corners[0] + desc.offset);
|
||||
f1 = kernel_data_fetch(attributes_float3, corners[1] + desc.offset);
|
||||
f2 = kernel_data_fetch(attributes_float3, corners[2] + desc.offset);
|
||||
f3 = kernel_data_fetch(attributes_float3, corners[3] + desc.offset);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -493,7 +491,7 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = make_float3(0.0f, 0.0f, 0.0f);
|
||||
|
||||
return kernel_tex_fetch(__attributes_float3, desc.offset);
|
||||
return kernel_data_fetch(attributes_float3, desc.offset);
|
||||
}
|
||||
else {
|
||||
if (dx)
|
||||
@@ -570,8 +568,7 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = zero_float4();
|
||||
|
||||
return kernel_tex_fetch(__attributes_float4,
|
||||
desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
return kernel_data_fetch(attributes_float4, desc.offset + subd_triangle_patch_face(kg, patch));
|
||||
}
|
||||
else if (desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
|
||||
float2 uv[3];
|
||||
@@ -579,10 +576,10 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
|
||||
|
||||
uint4 v = subd_triangle_patch_indices(kg, patch);
|
||||
|
||||
float4 f0 = kernel_tex_fetch(__attributes_float4, desc.offset + v.x);
|
||||
float4 f1 = kernel_tex_fetch(__attributes_float4, desc.offset + v.y);
|
||||
float4 f2 = kernel_tex_fetch(__attributes_float4, desc.offset + v.z);
|
||||
float4 f3 = kernel_tex_fetch(__attributes_float4, desc.offset + v.w);
|
||||
float4 f0 = kernel_data_fetch(attributes_float4, desc.offset + v.x);
|
||||
float4 f1 = kernel_data_fetch(attributes_float4, desc.offset + v.y);
|
||||
float4 f2 = kernel_data_fetch(attributes_float4, desc.offset + v.z);
|
||||
float4 f3 = kernel_data_fetch(attributes_float4, desc.offset + v.w);
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
f1 = (f1 + f0) * 0.5f;
|
||||
@@ -613,19 +610,19 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
|
||||
|
||||
if (desc.element == ATTR_ELEMENT_CORNER_BYTE) {
|
||||
f0 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, corners[0] + desc.offset)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, corners[0] + desc.offset)));
|
||||
f1 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, corners[1] + desc.offset)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, corners[1] + desc.offset)));
|
||||
f2 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, corners[2] + desc.offset)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, corners[2] + desc.offset)));
|
||||
f3 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, corners[3] + desc.offset)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, corners[3] + desc.offset)));
|
||||
}
|
||||
else {
|
||||
f0 = kernel_tex_fetch(__attributes_float4, corners[0] + desc.offset);
|
||||
f1 = kernel_tex_fetch(__attributes_float4, corners[1] + desc.offset);
|
||||
f2 = kernel_tex_fetch(__attributes_float4, corners[2] + desc.offset);
|
||||
f3 = kernel_tex_fetch(__attributes_float4, corners[3] + desc.offset);
|
||||
f0 = kernel_data_fetch(attributes_float4, corners[0] + desc.offset);
|
||||
f1 = kernel_data_fetch(attributes_float4, corners[1] + desc.offset);
|
||||
f2 = kernel_data_fetch(attributes_float4, corners[2] + desc.offset);
|
||||
f3 = kernel_data_fetch(attributes_float4, corners[3] + desc.offset);
|
||||
}
|
||||
|
||||
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
|
||||
@@ -652,7 +649,7 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
|
||||
if (dy)
|
||||
*dy = zero_float4();
|
||||
|
||||
return kernel_tex_fetch(__attributes_float4, desc.offset);
|
||||
return kernel_data_fetch(attributes_float4, desc.offset);
|
||||
}
|
||||
else {
|
||||
if (dx)
|
||||
|
@@ -15,10 +15,10 @@ CCL_NAMESPACE_BEGIN
|
||||
ccl_device_inline float3 triangle_normal(KernelGlobals kg, ccl_private ShaderData *sd)
|
||||
{
|
||||
/* load triangle vertices */
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
const float3 v0 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
|
||||
const float3 v1 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
|
||||
const float3 v2 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
const float3 v0 = kernel_data_fetch(tri_verts, tri_vindex.w + 0);
|
||||
const float3 v1 = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
|
||||
const float3 v2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
|
||||
|
||||
/* return normal */
|
||||
if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
|
||||
@@ -40,15 +40,15 @@ ccl_device_inline void triangle_point_normal(KernelGlobals kg,
|
||||
ccl_private int *shader)
|
||||
{
|
||||
/* load triangle vertices */
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
float3 v0 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
|
||||
float3 v1 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
|
||||
float3 v2 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
float3 v0 = kernel_data_fetch(tri_verts, tri_vindex.w + 0);
|
||||
float3 v1 = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
|
||||
float3 v2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
|
||||
/* compute point */
|
||||
float t = 1.0f - u - v;
|
||||
*P = (u * v0 + v * v1 + t * v2);
|
||||
/* get object flags */
|
||||
int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
int object_flag = kernel_data_fetch(object_flag, object);
|
||||
/* compute normal */
|
||||
if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
|
||||
*Ng = normalize(cross(v2 - v0, v1 - v0));
|
||||
@@ -57,17 +57,17 @@ ccl_device_inline void triangle_point_normal(KernelGlobals kg,
|
||||
*Ng = normalize(cross(v1 - v0, v2 - v0));
|
||||
}
|
||||
/* shader`*/
|
||||
*shader = kernel_tex_fetch(__tri_shader, prim);
|
||||
*shader = kernel_data_fetch(tri_shader, prim);
|
||||
}
|
||||
|
||||
/* Triangle vertex locations */
|
||||
|
||||
ccl_device_inline void triangle_vertices(KernelGlobals kg, int prim, float3 P[3])
|
||||
{
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
P[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
|
||||
P[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
|
||||
P[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
P[0] = kernel_data_fetch(tri_verts, tri_vindex.w + 0);
|
||||
P[1] = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
|
||||
P[2] = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
|
||||
}
|
||||
|
||||
/* Triangle vertex locations and vertex normals */
|
||||
@@ -77,13 +77,13 @@ ccl_device_inline void triangle_vertices_and_normals(KernelGlobals kg,
|
||||
float3 P[3],
|
||||
float3 N[3])
|
||||
{
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
P[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
|
||||
P[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
|
||||
P[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
|
||||
N[0] = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
|
||||
N[1] = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
|
||||
N[2] = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
P[0] = kernel_data_fetch(tri_verts, tri_vindex.w + 0);
|
||||
P[1] = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
|
||||
P[2] = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
|
||||
N[0] = kernel_data_fetch(tri_vnormal, tri_vindex.x);
|
||||
N[1] = kernel_data_fetch(tri_vnormal, tri_vindex.y);
|
||||
N[2] = kernel_data_fetch(tri_vnormal, tri_vindex.z);
|
||||
}
|
||||
|
||||
/* Interpolate smooth vertex normal from vertices */
|
||||
@@ -92,10 +92,10 @@ ccl_device_inline float3
|
||||
triangle_smooth_normal(KernelGlobals kg, float3 Ng, int prim, float u, float v)
|
||||
{
|
||||
/* load triangle vertices */
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
float3 n0 = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
|
||||
float3 n1 = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
|
||||
float3 n2 = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
float3 n0 = kernel_data_fetch(tri_vnormal, tri_vindex.x);
|
||||
float3 n1 = kernel_data_fetch(tri_vnormal, tri_vindex.y);
|
||||
float3 n2 = kernel_data_fetch(tri_vnormal, tri_vindex.z);
|
||||
|
||||
float3 N = safe_normalize((1.0f - u - v) * n2 + u * n0 + v * n1);
|
||||
|
||||
@@ -106,10 +106,10 @@ ccl_device_inline float3 triangle_smooth_normal_unnormalized(
|
||||
KernelGlobals kg, ccl_private const ShaderData *sd, float3 Ng, int prim, float u, float v)
|
||||
{
|
||||
/* load triangle vertices */
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
float3 n0 = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
|
||||
float3 n1 = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
|
||||
float3 n2 = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
float3 n0 = kernel_data_fetch(tri_vnormal, tri_vindex.x);
|
||||
float3 n1 = kernel_data_fetch(tri_vnormal, tri_vindex.y);
|
||||
float3 n2 = kernel_data_fetch(tri_vnormal, tri_vindex.z);
|
||||
|
||||
/* ensure that the normals are in object space */
|
||||
if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
|
||||
@@ -131,10 +131,10 @@ ccl_device_inline void triangle_dPdudv(KernelGlobals kg,
|
||||
ccl_private float3 *dPdv)
|
||||
{
|
||||
/* fetch triangle vertex coordinates */
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
|
||||
const float3 p0 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
|
||||
const float3 p1 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
|
||||
const float3 p2 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, prim);
|
||||
const float3 p0 = kernel_data_fetch(tri_verts, tri_vindex.w + 0);
|
||||
const float3 p1 = kernel_data_fetch(tri_verts, tri_vindex.w + 1);
|
||||
const float3 p2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2);
|
||||
|
||||
/* compute derivatives of P w.r.t. uv */
|
||||
*dPdu = (p0 - p2);
|
||||
@@ -153,16 +153,16 @@ ccl_device float triangle_attribute_float(KernelGlobals kg,
|
||||
float f0, f1, f2;
|
||||
|
||||
if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_VERTEX_MOTION)) {
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
f0 = kernel_tex_fetch(__attributes_float, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_tex_fetch(__attributes_float, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_tex_fetch(__attributes_float, desc.offset + tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
f0 = kernel_data_fetch(attributes_float, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_data_fetch(attributes_float, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_data_fetch(attributes_float, desc.offset + tri_vindex.z);
|
||||
}
|
||||
else {
|
||||
const int tri = desc.offset + sd->prim * 3;
|
||||
f0 = kernel_tex_fetch(__attributes_float, tri + 0);
|
||||
f1 = kernel_tex_fetch(__attributes_float, tri + 1);
|
||||
f2 = kernel_tex_fetch(__attributes_float, tri + 2);
|
||||
f0 = kernel_data_fetch(attributes_float, tri + 0);
|
||||
f1 = kernel_data_fetch(attributes_float, tri + 1);
|
||||
f2 = kernel_data_fetch(attributes_float, tri + 2);
|
||||
}
|
||||
|
||||
#ifdef __RAY_DIFFERENTIALS__
|
||||
@@ -185,7 +185,7 @@ ccl_device float triangle_attribute_float(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_FACE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_FACE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float, offset);
|
||||
return kernel_data_fetch(attributes_float, offset);
|
||||
}
|
||||
else {
|
||||
return 0.0f;
|
||||
@@ -203,16 +203,16 @@ ccl_device float2 triangle_attribute_float2(KernelGlobals kg,
|
||||
float2 f0, f1, f2;
|
||||
|
||||
if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_VERTEX_MOTION)) {
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
f0 = kernel_tex_fetch(__attributes_float2, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_tex_fetch(__attributes_float2, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_tex_fetch(__attributes_float2, desc.offset + tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
f0 = kernel_data_fetch(attributes_float2, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_data_fetch(attributes_float2, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_data_fetch(attributes_float2, desc.offset + tri_vindex.z);
|
||||
}
|
||||
else {
|
||||
const int tri = desc.offset + sd->prim * 3;
|
||||
f0 = kernel_tex_fetch(__attributes_float2, tri + 0);
|
||||
f1 = kernel_tex_fetch(__attributes_float2, tri + 1);
|
||||
f2 = kernel_tex_fetch(__attributes_float2, tri + 2);
|
||||
f0 = kernel_data_fetch(attributes_float2, tri + 0);
|
||||
f1 = kernel_data_fetch(attributes_float2, tri + 1);
|
||||
f2 = kernel_data_fetch(attributes_float2, tri + 2);
|
||||
}
|
||||
|
||||
#ifdef __RAY_DIFFERENTIALS__
|
||||
@@ -235,7 +235,7 @@ ccl_device float2 triangle_attribute_float2(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_FACE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_FACE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float2, offset);
|
||||
return kernel_data_fetch(attributes_float2, offset);
|
||||
}
|
||||
else {
|
||||
return make_float2(0.0f, 0.0f);
|
||||
@@ -253,16 +253,16 @@ ccl_device float3 triangle_attribute_float3(KernelGlobals kg,
|
||||
float3 f0, f1, f2;
|
||||
|
||||
if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_VERTEX_MOTION)) {
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
f0 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
f0 = kernel_data_fetch(attributes_float3, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_data_fetch(attributes_float3, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_data_fetch(attributes_float3, desc.offset + tri_vindex.z);
|
||||
}
|
||||
else {
|
||||
const int tri = desc.offset + sd->prim * 3;
|
||||
f0 = kernel_tex_fetch(__attributes_float3, tri + 0);
|
||||
f1 = kernel_tex_fetch(__attributes_float3, tri + 1);
|
||||
f2 = kernel_tex_fetch(__attributes_float3, tri + 2);
|
||||
f0 = kernel_data_fetch(attributes_float3, tri + 0);
|
||||
f1 = kernel_data_fetch(attributes_float3, tri + 1);
|
||||
f2 = kernel_data_fetch(attributes_float3, tri + 2);
|
||||
}
|
||||
|
||||
#ifdef __RAY_DIFFERENTIALS__
|
||||
@@ -285,7 +285,7 @@ ccl_device float3 triangle_attribute_float3(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_FACE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_FACE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float3, offset);
|
||||
return kernel_data_fetch(attributes_float3, offset);
|
||||
}
|
||||
else {
|
||||
return make_float3(0.0f, 0.0f, 0.0f);
|
||||
@@ -304,25 +304,25 @@ ccl_device float4 triangle_attribute_float4(KernelGlobals kg,
|
||||
float4 f0, f1, f2;
|
||||
|
||||
if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_VERTEX_MOTION)) {
|
||||
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
|
||||
f0 = kernel_tex_fetch(__attributes_float4, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_tex_fetch(__attributes_float4, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_tex_fetch(__attributes_float4, desc.offset + tri_vindex.z);
|
||||
const uint4 tri_vindex = kernel_data_fetch(tri_vindex, sd->prim);
|
||||
f0 = kernel_data_fetch(attributes_float4, desc.offset + tri_vindex.x);
|
||||
f1 = kernel_data_fetch(attributes_float4, desc.offset + tri_vindex.y);
|
||||
f2 = kernel_data_fetch(attributes_float4, desc.offset + tri_vindex.z);
|
||||
}
|
||||
else {
|
||||
const int tri = desc.offset + sd->prim * 3;
|
||||
if (desc.element == ATTR_ELEMENT_CORNER) {
|
||||
f0 = kernel_tex_fetch(__attributes_float4, tri + 0);
|
||||
f1 = kernel_tex_fetch(__attributes_float4, tri + 1);
|
||||
f2 = kernel_tex_fetch(__attributes_float4, tri + 2);
|
||||
f0 = kernel_data_fetch(attributes_float4, tri + 0);
|
||||
f1 = kernel_data_fetch(attributes_float4, tri + 1);
|
||||
f2 = kernel_data_fetch(attributes_float4, tri + 2);
|
||||
}
|
||||
else {
|
||||
f0 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, tri + 0)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, tri + 0)));
|
||||
f1 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, tri + 1)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, tri + 1)));
|
||||
f2 = color_srgb_to_linear_v4(
|
||||
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, tri + 2)));
|
||||
color_uchar4_to_float4(kernel_data_fetch(attributes_uchar4, tri + 2)));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -346,7 +346,7 @@ ccl_device float4 triangle_attribute_float4(KernelGlobals kg,
|
||||
if (desc.element & (ATTR_ELEMENT_FACE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
const int offset = (desc.element == ATTR_ELEMENT_FACE) ? desc.offset + sd->prim :
|
||||
desc.offset;
|
||||
return kernel_tex_fetch(__attributes_float4, offset);
|
||||
return kernel_data_fetch(attributes_float4, offset);
|
||||
}
|
||||
else {
|
||||
return zero_float4();
|
||||
|
@@ -23,17 +23,17 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg,
|
||||
int prim,
|
||||
int prim_addr)
|
||||
{
|
||||
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
|
||||
const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
|
||||
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
|
||||
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
|
||||
const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w;
|
||||
const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0),
|
||||
tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1),
|
||||
tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
|
||||
float t, u, v;
|
||||
if (ray_triangle_intersect(P, dir, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) {
|
||||
#ifdef __VISIBILITY_FLAG__
|
||||
/* Visibility flag test. we do it here under the assumption
|
||||
* that most triangles are culled by node flags.
|
||||
*/
|
||||
if (kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
|
||||
if (kernel_data_fetch(prim_visibility, prim_addr) & visibility)
|
||||
#endif
|
||||
{
|
||||
isect->object = object;
|
||||
@@ -66,10 +66,10 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
|
||||
ccl_private uint *lcg_state,
|
||||
int max_hits)
|
||||
{
|
||||
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
|
||||
const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
|
||||
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
|
||||
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
|
||||
const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w;
|
||||
const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0),
|
||||
tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1),
|
||||
tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
|
||||
float t, u, v;
|
||||
if (!ray_triangle_intersect(P, dir, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) {
|
||||
return false;
|
||||
@@ -139,10 +139,10 @@ ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg,
|
||||
const float u,
|
||||
const float v)
|
||||
{
|
||||
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
|
||||
const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
|
||||
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
|
||||
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
|
||||
const uint tri_vindex = kernel_data_fetch(tri_vindex, isect_prim).w;
|
||||
const packed_float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0),
|
||||
tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1),
|
||||
tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
|
||||
float w = 1.0f - u - v;
|
||||
|
||||
float3 P = u * tri_a + v * tri_b + w * tri_c;
|
||||
|
@@ -62,7 +62,7 @@ ccl_device float4 volume_attribute_float4(KernelGlobals kg,
|
||||
const AttributeDescriptor desc)
|
||||
{
|
||||
if (desc.element & (ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
|
||||
return kernel_tex_fetch(__attributes_float4, desc.offset);
|
||||
return kernel_data_fetch(attributes_float4, desc.offset);
|
||||
}
|
||||
else if (desc.element == ATTR_ELEMENT_VOXEL) {
|
||||
/* todo: optimize this so we don't have to transform both here and in
|
||||
|
@@ -160,7 +160,7 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
|
||||
int shader;
|
||||
triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
|
||||
|
||||
const int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
const int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
|
||||
Transform tfm = object_fetch_transform(kg, object, OBJECT_TRANSFORM);
|
||||
P = transform_point_auto(&tfm, P);
|
||||
@@ -193,7 +193,7 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
|
||||
}
|
||||
|
||||
const int shader_index = shader & SHADER_MASK;
|
||||
const int shader_flags = kernel_tex_fetch(__shaders, shader_index).flags;
|
||||
const int shader_flags = kernel_data_fetch(shaders, shader_index).flags;
|
||||
|
||||
/* Fast path for position and normal passes not affected by shaders. */
|
||||
if (kernel_data.film.pass_position != PASS_UNUSED) {
|
||||
|
@@ -122,7 +122,7 @@ ccl_device_forceinline void integrator_split_shadow_catcher(
|
||||
|
||||
/* Continue with shading shadow catcher surface. */
|
||||
const int shader = intersection_get_shader(kg, isect);
|
||||
const int flags = kernel_tex_fetch(__shaders, shader).flags;
|
||||
const int flags = kernel_data_fetch(shaders, shader).flags;
|
||||
const bool use_caustics = kernel_data.integrator.use_caustics &&
|
||||
(object_flags & SD_OBJECT_CAUSTICS);
|
||||
const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
|
||||
@@ -149,7 +149,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_shadow_catche
|
||||
integrator_state_read_isect(kg, state, &isect);
|
||||
|
||||
const int shader = intersection_get_shader(kg, &isect);
|
||||
const int flags = kernel_tex_fetch(__shaders, shader).flags;
|
||||
const int flags = kernel_data_fetch(shaders, shader).flags;
|
||||
const int object_flags = intersection_get_object_flags(kg, &isect);
|
||||
const bool use_caustics = kernel_data.integrator.use_caustics &&
|
||||
(object_flags & SD_OBJECT_CAUSTICS);
|
||||
@@ -203,7 +203,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
|
||||
if (!integrator_state_volume_stack_is_empty(kg, state)) {
|
||||
const bool hit_surface = hit && !(isect->type & PRIMITIVE_LAMP);
|
||||
const int shader = (hit_surface) ? intersection_get_shader(kg, isect) : SHADER_NONE;
|
||||
const int flags = (hit_surface) ? kernel_tex_fetch(__shaders, shader).flags : 0;
|
||||
const int flags = (hit_surface) ? kernel_data_fetch(shaders, shader).flags : 0;
|
||||
|
||||
if (!integrator_intersect_terminate(kg, state, flags)) {
|
||||
INTEGRATOR_PATH_NEXT(current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
|
||||
@@ -223,7 +223,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
|
||||
else {
|
||||
/* Hit a surface, continue with surface kernel unless terminated. */
|
||||
const int shader = intersection_get_shader(kg, isect);
|
||||
const int flags = kernel_tex_fetch(__shaders, shader).flags;
|
||||
const int flags = kernel_data_fetch(shaders, shader).flags;
|
||||
|
||||
if (!integrator_intersect_terminate(kg, state, flags)) {
|
||||
const int object_flags = intersection_get_object_flags(kg, isect);
|
||||
@@ -279,7 +279,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
|
||||
else {
|
||||
/* Hit a surface, continue with surface kernel unless terminated. */
|
||||
const int shader = intersection_get_shader(kg, isect);
|
||||
const int flags = kernel_tex_fetch(__shaders, shader).flags;
|
||||
const int flags = kernel_data_fetch(shaders, shader).flags;
|
||||
const int object_flags = intersection_get_object_flags(kg, isect);
|
||||
const bool use_caustics = kernel_data.integrator.use_caustics &&
|
||||
(object_flags & SD_OBJECT_CAUSTICS);
|
||||
@@ -332,7 +332,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
|
||||
ray.t = kernel_data.integrator.ao_bounces_distance;
|
||||
|
||||
if (last_isect_object != OBJECT_NONE) {
|
||||
const float object_ao_distance = kernel_tex_fetch(__objects, last_isect_object).ao_distance;
|
||||
const float object_ao_distance = kernel_data_fetch(objects, last_isect_object).ao_distance;
|
||||
if (object_ao_distance != 0.0f) {
|
||||
ray.t = object_ao_distance;
|
||||
}
|
||||
@@ -366,7 +366,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
|
||||
bool from_caustic_caster = false;
|
||||
bool from_caustic_receiver = false;
|
||||
if (!(path_flag & PATH_RAY_CAMERA) && last_isect_object != OBJECT_NONE) {
|
||||
const int object_flags = kernel_tex_fetch(__object_flag, last_isect_object);
|
||||
const int object_flags = kernel_data_fetch(object_flag, last_isect_object);
|
||||
from_caustic_receiver = (object_flags & SD_OBJECT_CAUSTICS_RECEIVER);
|
||||
from_caustic_caster = (object_flags & SD_OBJECT_CAUSTICS_CASTER);
|
||||
}
|
||||
|
@@ -115,7 +115,7 @@ ccl_device_forceinline void mnee_update_light_sample(KernelGlobals kg,
|
||||
{
|
||||
/* correct light sample position/direction and pdf
|
||||
* NOTE: preserve pdf in area measure */
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, ls->lamp);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, ls->lamp);
|
||||
|
||||
if (ls->type == LIGHT_POINT || ls->type == LIGHT_SPOT) {
|
||||
ls->D = normalize_len(ls->P - P, &ls->t);
|
||||
@@ -154,12 +154,12 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg,
|
||||
ccl_private const Intersection *isect,
|
||||
ccl_private ShaderData *sd_vtx)
|
||||
{
|
||||
sd_vtx->object = (isect->object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, isect->prim) :
|
||||
sd_vtx->object = (isect->object == OBJECT_NONE) ? kernel_data_fetch(prim_object, isect->prim) :
|
||||
isect->object;
|
||||
|
||||
sd_vtx->type = isect->type;
|
||||
sd_vtx->flag = 0;
|
||||
sd_vtx->object_flag = kernel_tex_fetch(__object_flag, sd_vtx->object);
|
||||
sd_vtx->object_flag = kernel_data_fetch(object_flag, sd_vtx->object);
|
||||
|
||||
/* Matrices and time. */
|
||||
shader_setup_object_transforms(kg, sd_vtx, ray->time);
|
||||
@@ -171,7 +171,7 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg,
|
||||
sd_vtx->u = isect->u;
|
||||
sd_vtx->v = isect->v;
|
||||
|
||||
sd_vtx->shader = kernel_tex_fetch(__tri_shader, sd_vtx->prim);
|
||||
sd_vtx->shader = kernel_data_fetch(tri_shader, sd_vtx->prim);
|
||||
|
||||
float3 verts[3];
|
||||
float3 normals[3];
|
||||
@@ -509,7 +509,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
|
||||
break;
|
||||
|
||||
int hit_object = (projection_isect.object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, projection_isect.prim) :
|
||||
kernel_data_fetch(prim_object, projection_isect.prim) :
|
||||
projection_isect.object;
|
||||
|
||||
if (hit_object == mv.object) {
|
||||
@@ -870,7 +870,7 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg,
|
||||
probe_ray.D = normalize_len(v.p - probe_ray.P, &probe_ray.t);
|
||||
if (scene_intersect(kg, &probe_ray, PATH_RAY_TRANSMIT, &probe_isect)) {
|
||||
int hit_object = (probe_isect.object == OBJECT_NONE) ?
|
||||
kernel_tex_fetch(__prim_object, probe_isect.prim) :
|
||||
kernel_data_fetch(prim_object, probe_isect.prim) :
|
||||
probe_isect.object;
|
||||
/* Test whether the ray hit the appropriate object at its intended location. */
|
||||
if (hit_object != v.object || fabsf(probe_ray.t - probe_isect.t) > MNEE_MIN_DISTANCE)
|
||||
|
@@ -107,7 +107,7 @@ ccl_device_inline void integrate_background(KernelGlobals kg,
|
||||
for (int lamp = 0; lamp < kernel_data.integrator.num_all_lights; lamp++) {
|
||||
/* This path should have been resolved with mnee, it will
|
||||
* generate a firefly for small lights since it is improbable. */
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, lamp);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
|
||||
if (klight->type == LIGHT_BACKGROUND && klight->use_caustics) {
|
||||
eval_background = false;
|
||||
break;
|
||||
@@ -160,7 +160,7 @@ ccl_device_inline void integrate_distant_lights(KernelGlobals kg,
|
||||
if (INTEGRATOR_STATE(state, path, mnee) & PATH_MNEE_CULL_LIGHT_CONNECTION) {
|
||||
/* This path should have been resolved with mnee, it will
|
||||
* generate a firefly for small lights since it is improbable. */
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, lamp);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
|
||||
if (klight->use_caustics)
|
||||
return;
|
||||
}
|
||||
|
@@ -141,7 +141,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
|
||||
{
|
||||
if (ls.lamp != LAMP_NONE) {
|
||||
/* Is this a caustic light? */
|
||||
const bool use_caustics = kernel_tex_fetch(__lights, ls.lamp).use_caustics;
|
||||
const bool use_caustics = kernel_data_fetch(lights, ls.lamp).use_caustics;
|
||||
if (use_caustics) {
|
||||
/* Are we on a caustic caster? */
|
||||
if (is_transmission && (sd->object_flag & SD_OBJECT_CAUSTICS_CASTER))
|
||||
|
@@ -528,12 +528,12 @@ ccl_device bool shader_constant_emission_eval(KernelGlobals kg,
|
||||
ccl_private float3 *eval)
|
||||
{
|
||||
int shader_index = shader & SHADER_MASK;
|
||||
int shader_flag = kernel_tex_fetch(__shaders, shader_index).flags;
|
||||
int shader_flag = kernel_data_fetch(shaders, shader_index).flags;
|
||||
|
||||
if (shader_flag & SD_HAS_CONSTANT_EMISSION) {
|
||||
*eval = make_float3(kernel_tex_fetch(__shaders, shader_index).constant_emission[0],
|
||||
kernel_tex_fetch(__shaders, shader_index).constant_emission[1],
|
||||
kernel_tex_fetch(__shaders, shader_index).constant_emission[2]);
|
||||
*eval = make_float3(kernel_data_fetch(shaders, shader_index).constant_emission[0],
|
||||
kernel_data_fetch(shaders, shader_index).constant_emission[1],
|
||||
kernel_data_fetch(shaders, shader_index).constant_emission[2]);
|
||||
|
||||
return true;
|
||||
}
|
||||
@@ -821,11 +821,11 @@ ccl_device_inline void shader_eval_volume(KernelGlobals kg,
|
||||
sd->shader = entry.shader;
|
||||
|
||||
sd->flag &= ~SD_SHADER_FLAGS;
|
||||
sd->flag |= kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->flag |= kernel_data_fetch(shaders, (sd->shader & SHADER_MASK)).flags;
|
||||
sd->object_flag &= ~SD_OBJECT_FLAGS;
|
||||
|
||||
if (sd->object != OBJECT_NONE) {
|
||||
sd->object_flag |= kernel_tex_fetch(__object_flag, sd->object);
|
||||
sd->object_flag |= kernel_data_fetch(object_flag, sd->object);
|
||||
|
||||
# ifdef __OBJECT_MOTION__
|
||||
/* todo: this is inefficient for motion blur, we should be
|
||||
@@ -837,7 +837,7 @@ ccl_device_inline void shader_eval_volume(KernelGlobals kg,
|
||||
kernel_assert(v_desc.offset != ATTR_STD_NOT_FOUND);
|
||||
|
||||
const float3 P = sd->P;
|
||||
const float velocity_scale = kernel_tex_fetch(__objects, sd->object).velocity_scale;
|
||||
const float velocity_scale = kernel_data_fetch(objects, sd->object).velocity_scale;
|
||||
const float time_offset = kernel_data.cam.motion_position == MOTION_POSITION_CENTER ?
|
||||
0.5f :
|
||||
0.0f;
|
||||
@@ -946,7 +946,7 @@ ccl_device void shader_eval_displacement(KernelGlobals kg,
|
||||
|
||||
ccl_device float shader_cryptomatte_id(KernelGlobals kg, int shader)
|
||||
{
|
||||
return kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).cryptomatte_id;
|
||||
return kernel_data_fetch(shaders, (shader & SHADER_MASK)).cryptomatte_id;
|
||||
}
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
@@ -147,7 +147,7 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
|
||||
/* Update volume stack if needed. */
|
||||
if (kernel_data.integrator.use_volumes) {
|
||||
const int object = ss_isect.hits[0].object;
|
||||
const int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
const int object_flag = kernel_data_fetch(object_flag, object);
|
||||
|
||||
if (object_flag & SD_OBJECT_INTERSECTS_VOLUME) {
|
||||
float3 P = INTEGRATOR_STATE(state, ray, P);
|
||||
@@ -170,7 +170,7 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
|
||||
INTEGRATOR_STATE_WRITE(state, path, rng_offset) += PRNG_BOUNCE_NUM;
|
||||
|
||||
const int shader = intersection_get_shader(kg, &ss_isect.hits[0]);
|
||||
const int shader_flags = kernel_tex_fetch(__shaders, shader).flags;
|
||||
const int shader_flags = kernel_data_fetch(shaders, shader).flags;
|
||||
const int object_flags = intersection_get_object_flags(kg, &ss_isect.hits[0]);
|
||||
const bool use_caustics = kernel_data.integrator.use_caustics &&
|
||||
(object_flags & SD_OBJECT_CAUSTICS);
|
||||
|
@@ -113,7 +113,7 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
|
||||
for (int hit = 0; hit < num_eval_hits; hit++) {
|
||||
/* Get geometric normal. */
|
||||
const int object = ss_isect.hits[hit].object;
|
||||
const int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
const int object_flag = kernel_data_fetch(object_flag, object);
|
||||
float3 hit_Ng = ss_isect.Ng[hit];
|
||||
if (path_flag & PATH_RAY_SUBSURFACE_BACKFACING) {
|
||||
hit_Ng = -hit_Ng;
|
||||
|
@@ -133,7 +133,7 @@ ccl_device float volume_stack_step_size(KernelGlobals kg, StackReadOp stack_read
|
||||
break;
|
||||
}
|
||||
|
||||
int shader_flag = kernel_tex_fetch(__shaders, (entry.shader & SHADER_MASK)).flags;
|
||||
int shader_flag = kernel_data_fetch(shaders, (entry.shader & SHADER_MASK)).flags;
|
||||
|
||||
bool heterogeneous = false;
|
||||
|
||||
@@ -146,7 +146,7 @@ ccl_device float volume_stack_step_size(KernelGlobals kg, StackReadOp stack_read
|
||||
* heterogeneous volume objects may be using the same shader. */
|
||||
int object = entry.object;
|
||||
if (object != OBJECT_NONE) {
|
||||
int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (object_flag & SD_OBJECT_HAS_VOLUME_ATTRIBUTES) {
|
||||
heterogeneous = true;
|
||||
}
|
||||
@@ -180,7 +180,7 @@ ccl_device VolumeSampleMethod volume_stack_sample_method(KernelGlobals kg, Integ
|
||||
break;
|
||||
}
|
||||
|
||||
int shader_flag = kernel_tex_fetch(__shaders, (entry.shader & SHADER_MASK)).flags;
|
||||
int shader_flag = kernel_data_fetch(shaders, (entry.shader & SHADER_MASK)).flags;
|
||||
|
||||
if (shader_flag & SD_VOLUME_MIS) {
|
||||
/* Multiple importance sampling. */
|
||||
|
@@ -31,7 +31,7 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
|
||||
int step = count >> 1;
|
||||
int middle = first + step;
|
||||
|
||||
if (kernel_tex_fetch(__light_background_marginal_cdf, middle).y < randv) {
|
||||
if (kernel_data_fetch(light_background_marginal_cdf, middle).y < randv) {
|
||||
first = middle + 1;
|
||||
count -= step + 1;
|
||||
}
|
||||
@@ -42,9 +42,9 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
|
||||
int index_v = max(0, first - 1);
|
||||
kernel_assert(index_v >= 0 && index_v < res_y);
|
||||
|
||||
float2 cdf_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v);
|
||||
float2 cdf_next_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v + 1);
|
||||
float2 cdf_last_v = kernel_tex_fetch(__light_background_marginal_cdf, res_y);
|
||||
float2 cdf_v = kernel_data_fetch(light_background_marginal_cdf, index_v);
|
||||
float2 cdf_next_v = kernel_data_fetch(light_background_marginal_cdf, index_v + 1);
|
||||
float2 cdf_last_v = kernel_data_fetch(light_background_marginal_cdf, res_y);
|
||||
|
||||
/* importance-sampled V direction */
|
||||
float dv = inverse_lerp(cdf_v.y, cdf_next_v.y, randv);
|
||||
@@ -57,7 +57,7 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
|
||||
int step = count >> 1;
|
||||
int middle = first + step;
|
||||
|
||||
if (kernel_tex_fetch(__light_background_conditional_cdf, index_v * cdf_width + middle).y <
|
||||
if (kernel_data_fetch(light_background_conditional_cdf, index_v * cdf_width + middle).y <
|
||||
randu) {
|
||||
first = middle + 1;
|
||||
count -= step + 1;
|
||||
@@ -69,12 +69,12 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
|
||||
int index_u = max(0, first - 1);
|
||||
kernel_assert(index_u >= 0 && index_u < res_x);
|
||||
|
||||
float2 cdf_u = kernel_tex_fetch(__light_background_conditional_cdf,
|
||||
index_v * cdf_width + index_u);
|
||||
float2 cdf_next_u = kernel_tex_fetch(__light_background_conditional_cdf,
|
||||
index_v * cdf_width + index_u + 1);
|
||||
float2 cdf_last_u = kernel_tex_fetch(__light_background_conditional_cdf,
|
||||
index_v * cdf_width + res_x);
|
||||
float2 cdf_u = kernel_data_fetch(light_background_conditional_cdf,
|
||||
index_v * cdf_width + index_u);
|
||||
float2 cdf_next_u = kernel_data_fetch(light_background_conditional_cdf,
|
||||
index_v * cdf_width + index_u + 1);
|
||||
float2 cdf_last_u = kernel_data_fetch(light_background_conditional_cdf,
|
||||
index_v * cdf_width + res_x);
|
||||
|
||||
/* importance-sampled U direction */
|
||||
float du = inverse_lerp(cdf_u.y, cdf_next_u.y, randu);
|
||||
@@ -112,9 +112,9 @@ ccl_device float background_map_pdf(KernelGlobals kg, float3 direction)
|
||||
int index_v = clamp(float_to_int(uv.y * res_y), 0, res_y - 1);
|
||||
|
||||
/* pdfs in V direction */
|
||||
float2 cdf_last_u = kernel_tex_fetch(__light_background_conditional_cdf,
|
||||
index_v * cdf_width + res_x);
|
||||
float2 cdf_last_v = kernel_tex_fetch(__light_background_marginal_cdf, res_y);
|
||||
float2 cdf_last_u = kernel_data_fetch(light_background_conditional_cdf,
|
||||
index_v * cdf_width + res_x);
|
||||
float2 cdf_last_v = kernel_data_fetch(light_background_marginal_cdf, res_y);
|
||||
|
||||
float denom = (M_2PI_F * M_PI_F * sin_theta) * cdf_last_u.x * cdf_last_v.x;
|
||||
|
||||
@@ -122,9 +122,9 @@ ccl_device float background_map_pdf(KernelGlobals kg, float3 direction)
|
||||
return 0.0f;
|
||||
|
||||
/* pdfs in U direction */
|
||||
float2 cdf_u = kernel_tex_fetch(__light_background_conditional_cdf,
|
||||
index_v * cdf_width + index_u);
|
||||
float2 cdf_v = kernel_tex_fetch(__light_background_marginal_cdf, index_v);
|
||||
float2 cdf_u = kernel_data_fetch(light_background_conditional_cdf,
|
||||
index_v * cdf_width + index_u);
|
||||
float2 cdf_v = kernel_data_fetch(light_background_marginal_cdf, index_v);
|
||||
|
||||
return (cdf_u.x * cdf_v.x) / denom;
|
||||
}
|
||||
@@ -133,7 +133,7 @@ ccl_device_inline bool background_portal_data_fetch_and_check_side(
|
||||
KernelGlobals kg, float3 P, int index, ccl_private float3 *lightpos, ccl_private float3 *dir)
|
||||
{
|
||||
int portal = kernel_data.background.portal_offset + index;
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, portal);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, portal);
|
||||
|
||||
*lightpos = make_float3(klight->co[0], klight->co[1], klight->co[2]);
|
||||
*dir = make_float3(klight->area.dir[0], klight->area.dir[1], klight->area.dir[2]);
|
||||
@@ -166,7 +166,7 @@ ccl_device_inline float background_portal_pdf(
|
||||
num_possible++;
|
||||
|
||||
int portal = kernel_data.background.portal_offset + p;
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, portal);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, portal);
|
||||
float3 axisu = make_float3(
|
||||
klight->area.axisu[0], klight->area.axisu[1], klight->area.axisu[2]);
|
||||
float3 axisv = make_float3(
|
||||
@@ -242,7 +242,7 @@ ccl_device float3 background_portal_sample(KernelGlobals kg,
|
||||
if (portal == 0) {
|
||||
/* p is the portal to be sampled. */
|
||||
int portal = kernel_data.background.portal_offset + p;
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, portal);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, portal);
|
||||
float3 axisu = make_float3(
|
||||
klight->area.axisu[0], klight->area.axisu[1], klight->area.axisu[2]);
|
||||
float3 axisv = make_float3(
|
||||
|
@@ -38,7 +38,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
|
||||
const uint32_t path_flag,
|
||||
ccl_private LightSample *ls)
|
||||
{
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, lamp);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
|
||||
if (path_flag & PATH_RAY_SHADOW_CATCHER_PASS) {
|
||||
if (klight->shader_id & SHADER_EXCLUDE_SHADOW_CATCHER) {
|
||||
return false;
|
||||
@@ -237,7 +237,7 @@ ccl_device bool lights_intersect(KernelGlobals kg,
|
||||
const uint32_t path_flag)
|
||||
{
|
||||
for (int lamp = 0; lamp < kernel_data.integrator.num_all_lights; lamp++) {
|
||||
const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, lamp);
|
||||
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
|
||||
|
||||
if (path_flag & PATH_RAY_CAMERA) {
|
||||
if (klight->shader_id & SHADER_EXCLUDE_CAMERA) {
|
||||
@@ -358,7 +358,7 @@ ccl_device bool light_sample_from_distant_ray(KernelGlobals kg,
|
||||
const int lamp,
|
||||
ccl_private LightSample *ccl_restrict ls)
|
||||
{
|
||||
ccl_global const KernelLight *klight = &kernel_tex_fetch(__lights, lamp);
|
||||
ccl_global const KernelLight *klight = &kernel_data_fetch(lights, lamp);
|
||||
const int shader = klight->shader_id;
|
||||
const float radius = klight->distant.radius;
|
||||
const LightType type = (LightType)klight->type;
|
||||
@@ -433,7 +433,7 @@ ccl_device bool light_sample_from_intersection(KernelGlobals kg,
|
||||
ccl_private LightSample *ccl_restrict ls)
|
||||
{
|
||||
const int lamp = isect->prim;
|
||||
ccl_global const KernelLight *klight = &kernel_tex_fetch(__lights, lamp);
|
||||
ccl_global const KernelLight *klight = &kernel_data_fetch(lights, lamp);
|
||||
LightType type = (LightType)klight->type;
|
||||
ls->type = type;
|
||||
ls->shader = klight->shader_id;
|
||||
@@ -562,7 +562,7 @@ ccl_device_inline bool triangle_world_space_vertices(
|
||||
KernelGlobals kg, int object, int prim, float time, float3 V[3])
|
||||
{
|
||||
bool has_motion = false;
|
||||
const int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
const int object_flag = kernel_data_fetch(object_flag, object);
|
||||
|
||||
if (object_flag & SD_OBJECT_HAS_VERTEX_MOTION && time >= 0.0f) {
|
||||
motion_triangle_vertices(kg, object, prim, time, V);
|
||||
@@ -699,12 +699,12 @@ ccl_device_forceinline void triangle_light_sample(KernelGlobals kg,
|
||||
float area = 0.5f * Nl;
|
||||
|
||||
/* flip normal if necessary */
|
||||
const int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
const int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
|
||||
ls->Ng = -ls->Ng;
|
||||
}
|
||||
ls->eval_fac = 1.0f;
|
||||
ls->shader = kernel_tex_fetch(__tri_shader, prim);
|
||||
ls->shader = kernel_data_fetch(tri_shader, prim);
|
||||
ls->object = object;
|
||||
ls->prim = prim;
|
||||
ls->lamp = LAMP_NONE;
|
||||
@@ -845,7 +845,7 @@ ccl_device int light_distribution_sample(KernelGlobals kg, ccl_private float *ra
|
||||
int half_len = len >> 1;
|
||||
int middle = first + half_len;
|
||||
|
||||
if (r < kernel_tex_fetch(__light_distribution, middle).totarea) {
|
||||
if (r < kernel_data_fetch(light_distribution, middle).totarea) {
|
||||
len = half_len;
|
||||
}
|
||||
else {
|
||||
@@ -860,8 +860,8 @@ ccl_device int light_distribution_sample(KernelGlobals kg, ccl_private float *ra
|
||||
|
||||
/* Rescale to reuse random number. this helps the 2D samples within
|
||||
* each area light be stratified as well. */
|
||||
float distr_min = kernel_tex_fetch(__light_distribution, index).totarea;
|
||||
float distr_max = kernel_tex_fetch(__light_distribution, index + 1).totarea;
|
||||
float distr_min = kernel_data_fetch(light_distribution, index).totarea;
|
||||
float distr_max = kernel_data_fetch(light_distribution, index + 1).totarea;
|
||||
*randu = (r - distr_min) / (distr_max - distr_min);
|
||||
|
||||
return index;
|
||||
@@ -871,7 +871,7 @@ ccl_device int light_distribution_sample(KernelGlobals kg, ccl_private float *ra
|
||||
|
||||
ccl_device_inline bool light_select_reached_max_bounces(KernelGlobals kg, int index, int bounce)
|
||||
{
|
||||
return (bounce > kernel_tex_fetch(__lights, index).max_bounces);
|
||||
return (bounce > kernel_data_fetch(lights, index).max_bounces);
|
||||
}
|
||||
|
||||
template<bool in_volume_segment>
|
||||
@@ -886,8 +886,8 @@ ccl_device_noinline bool light_distribution_sample(KernelGlobals kg,
|
||||
{
|
||||
/* Sample light index from distribution. */
|
||||
const int index = light_distribution_sample(kg, &randu);
|
||||
ccl_global const KernelLightDistribution *kdistribution = &kernel_tex_fetch(__light_distribution,
|
||||
index);
|
||||
ccl_global const KernelLightDistribution *kdistribution = &kernel_data_fetch(light_distribution,
|
||||
index);
|
||||
const int prim = kdistribution->prim;
|
||||
|
||||
if (prim >= 0) {
|
||||
@@ -896,7 +896,7 @@ ccl_device_noinline bool light_distribution_sample(KernelGlobals kg,
|
||||
|
||||
/* Exclude synthetic meshes from shadow catcher pass. */
|
||||
if ((path_flag & PATH_RAY_SHADOW_CATCHER_PASS) &&
|
||||
!(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_SHADOW_CATCHER)) {
|
||||
!(kernel_data_fetch(object_flag, object) & SD_OBJECT_SHADOW_CATCHER)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
|
@@ -81,7 +81,7 @@ light_sample_shader_eval(KernelGlobals kg,
|
||||
eval *= ls->eval_fac;
|
||||
|
||||
if (ls->lamp != LAMP_NONE) {
|
||||
ccl_global const KernelLight *klight = &kernel_tex_fetch(__lights, ls->lamp);
|
||||
ccl_global const KernelLight *klight = &kernel_data_fetch(lights, ls->lamp);
|
||||
eval *= make_float3(klight->strength[0], klight->strength[1], klight->strength[2]);
|
||||
}
|
||||
|
||||
@@ -187,7 +187,7 @@ ccl_device_inline float3 shadow_ray_offset(KernelGlobals kg,
|
||||
|
||||
if ((sd->type & PRIMITIVE_TRIANGLE) && (sd->shader & SHADER_SMOOTH_NORMAL)) {
|
||||
const float offset_cutoff =
|
||||
kernel_tex_fetch(__objects, sd->object).shadow_terminator_geometry_offset;
|
||||
kernel_data_fetch(objects, sd->object).shadow_terminator_geometry_offset;
|
||||
/* Do ray offset (heavy stuff) only for close to be terminated triangles:
|
||||
* offset_cutoff = 0.1f means that 10-20% of rays will be affected. Also
|
||||
* make a smooth transition near the threshold. */
|
||||
|
@@ -97,7 +97,7 @@ ccl_device float pmj_sample_1D(KernelGlobals kg, uint sample, uint rng_hash, uin
|
||||
* the x part is used for even dims and the y for odd. */
|
||||
int index = 2 * ((dim >> 1) * NUM_PMJ_SAMPLES + (s % NUM_PMJ_SAMPLES)) + (dim & 1);
|
||||
|
||||
float fx = kernel_tex_fetch(__sample_pattern_lut, index);
|
||||
float fx = kernel_data_fetch(sample_pattern_lut, index);
|
||||
|
||||
#ifndef _NO_CRANLEY_PATTERSON_ROTATION_
|
||||
/* Use Cranley-Patterson rotation to displace the sample pattern. */
|
||||
@@ -154,8 +154,8 @@ ccl_device void pmj_sample_2D(KernelGlobals kg,
|
||||
uint dim = d % NUM_PMJ_PATTERNS;
|
||||
int index = 2 * (dim * NUM_PMJ_SAMPLES + (s % NUM_PMJ_SAMPLES));
|
||||
|
||||
float fx = kernel_tex_fetch(__sample_pattern_lut, index);
|
||||
float fy = kernel_tex_fetch(__sample_pattern_lut, index + 1);
|
||||
float fx = kernel_data_fetch(sample_pattern_lut, index);
|
||||
float fy = kernel_data_fetch(sample_pattern_lut, index + 1);
|
||||
|
||||
#ifndef _NO_CRANLEY_PATTERSON_ROTATION_
|
||||
/* Use Cranley-Patterson rotation to displace the sample pattern. */
|
||||
|
@@ -32,7 +32,7 @@ ccl_device uint sobol_dimension(KernelGlobals kg, int index, int dimension)
|
||||
uint i = index + SOBOL_SKIP;
|
||||
for (int j = 0, x; (x = find_first_set(i)); i >>= x) {
|
||||
j += x;
|
||||
result ^= __float_as_uint(kernel_tex_fetch(__sample_pattern_lut, 32 * dimension + j - 1));
|
||||
result ^= __float_as_uint(kernel_data_fetch(sample_pattern_lut, 32 * dimension + j - 1));
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
@@ -222,7 +222,7 @@ ccl_device float3 svm_bevel(
|
||||
/* Get geometric normal. */
|
||||
float3 hit_Ng = isect.Ng[hit];
|
||||
int object = isect.hits[hit].object;
|
||||
int object_flag = kernel_tex_fetch(__object_flag, object);
|
||||
int object_flag = kernel_data_fetch(object_flag, object);
|
||||
if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
|
||||
hit_Ng = -hit_Ng;
|
||||
}
|
||||
@@ -230,7 +230,7 @@ ccl_device float3 svm_bevel(
|
||||
/* Compute smooth normal. */
|
||||
float3 N = hit_Ng;
|
||||
int prim = isect.hits[hit].prim;
|
||||
int shader = kernel_tex_fetch(__tri_shader, prim);
|
||||
int shader = kernel_data_fetch(tri_shader, prim);
|
||||
|
||||
if (shader & SHADER_SMOOTH_NORMAL) {
|
||||
float u = isect.hits[hit].u;
|
||||
|
@@ -17,7 +17,7 @@ ccl_device_inline float interpolate_ies_vertical(
|
||||
* Therefore, the assumption is made that the light is going to be symmetrical, which means that
|
||||
* we can just take the corresponding value at the current horizontal coordinate. */
|
||||
|
||||
#define IES_LOOKUP(v) kernel_tex_fetch(__ies, ofs + h * v_num + (v))
|
||||
#define IES_LOOKUP(v) kernel_data_fetch(ies, ofs + h * v_num + (v))
|
||||
/* If v is zero, assume symmetry and read at v=1 instead of v=-1. */
|
||||
float a = IES_LOOKUP((v == 0) ? 1 : v - 1);
|
||||
float b = IES_LOOKUP(v);
|
||||
@@ -31,16 +31,16 @@ ccl_device_inline float interpolate_ies_vertical(
|
||||
ccl_device_inline float kernel_ies_interp(KernelGlobals kg, int slot, float h_angle, float v_angle)
|
||||
{
|
||||
/* Find offset of the IES data in the table. */
|
||||
int ofs = __float_as_int(kernel_tex_fetch(__ies, slot));
|
||||
int ofs = __float_as_int(kernel_data_fetch(ies, slot));
|
||||
if (ofs == -1) {
|
||||
return 100.0f;
|
||||
}
|
||||
|
||||
int h_num = __float_as_int(kernel_tex_fetch(__ies, ofs++));
|
||||
int v_num = __float_as_int(kernel_tex_fetch(__ies, ofs++));
|
||||
int h_num = __float_as_int(kernel_data_fetch(ies, ofs++));
|
||||
int v_num = __float_as_int(kernel_data_fetch(ies, ofs++));
|
||||
|
||||
#define IES_LOOKUP_ANGLE_H(h) kernel_tex_fetch(__ies, ofs + (h))
|
||||
#define IES_LOOKUP_ANGLE_V(v) kernel_tex_fetch(__ies, ofs + h_num + (v))
|
||||
#define IES_LOOKUP_ANGLE_H(h) kernel_data_fetch(ies, ofs + (h))
|
||||
#define IES_LOOKUP_ANGLE_V(v) kernel_data_fetch(ies, ofs + h_num + (v))
|
||||
|
||||
/* Check whether the angle is within the bounds of the IES texture. */
|
||||
if (v_angle >= IES_LOOKUP_ANGLE_V(v_num - 1)) {
|
||||
|
@@ -9,7 +9,7 @@ CCL_NAMESPACE_BEGIN
|
||||
|
||||
ccl_device_inline float fetch_float(KernelGlobals kg, int offset)
|
||||
{
|
||||
uint4 node = kernel_tex_fetch(__svm_nodes, offset);
|
||||
uint4 node = kernel_data_fetch(svm_nodes, offset);
|
||||
return __uint_as_float(node.x);
|
||||
}
|
||||
|
||||
|
@@ -95,14 +95,14 @@ ccl_device_inline bool stack_valid(uint a)
|
||||
|
||||
ccl_device_inline uint4 read_node(KernelGlobals kg, ccl_private int *offset)
|
||||
{
|
||||
uint4 node = kernel_tex_fetch(__svm_nodes, *offset);
|
||||
uint4 node = kernel_data_fetch(svm_nodes, *offset);
|
||||
(*offset)++;
|
||||
return node;
|
||||
}
|
||||
|
||||
ccl_device_inline float4 read_node_float(KernelGlobals kg, ccl_private int *offset)
|
||||
{
|
||||
uint4 node = kernel_tex_fetch(__svm_nodes, *offset);
|
||||
uint4 node = kernel_data_fetch(svm_nodes, *offset);
|
||||
float4 f = make_float4(__uint_as_float(node.x),
|
||||
__uint_as_float(node.y),
|
||||
__uint_as_float(node.z),
|
||||
@@ -113,7 +113,7 @@ ccl_device_inline float4 read_node_float(KernelGlobals kg, ccl_private int *offs
|
||||
|
||||
ccl_device_inline float4 fetch_node_float(KernelGlobals kg, int offset)
|
||||
{
|
||||
uint4 node = kernel_tex_fetch(__svm_nodes, offset);
|
||||
uint4 node = kernel_data_fetch(svm_nodes, offset);
|
||||
return make_float4(__uint_as_float(node.x),
|
||||
__uint_as_float(node.y),
|
||||
__uint_as_float(node.z),
|
||||
|
@@ -1,82 +0,0 @@
|
||||
/* SPDX-License-Identifier: Apache-2.0
|
||||
* Copyright 2011-2022 Blender Foundation */
|
||||
|
||||
#ifndef KERNEL_TEX
|
||||
# define KERNEL_TEX(type, name)
|
||||
#endif
|
||||
|
||||
/* BVH2, not used for OptiX or Embree. */
|
||||
KERNEL_TEX(float4, __bvh_nodes)
|
||||
KERNEL_TEX(float4, __bvh_leaf_nodes)
|
||||
KERNEL_TEX(uint, __prim_type)
|
||||
KERNEL_TEX(uint, __prim_visibility)
|
||||
KERNEL_TEX(uint, __prim_index)
|
||||
KERNEL_TEX(uint, __prim_object)
|
||||
KERNEL_TEX(uint, __object_node)
|
||||
KERNEL_TEX(float2, __prim_time)
|
||||
|
||||
/* objects */
|
||||
KERNEL_TEX(KernelObject, __objects)
|
||||
KERNEL_TEX(Transform, __object_motion_pass)
|
||||
KERNEL_TEX(DecomposedTransform, __object_motion)
|
||||
KERNEL_TEX(uint, __object_flag)
|
||||
KERNEL_TEX(float, __object_volume_step)
|
||||
KERNEL_TEX(uint, __object_prim_offset)
|
||||
|
||||
/* cameras */
|
||||
KERNEL_TEX(DecomposedTransform, __camera_motion)
|
||||
|
||||
/* triangles */
|
||||
KERNEL_TEX(uint, __tri_shader)
|
||||
KERNEL_TEX(packed_float3, __tri_vnormal)
|
||||
KERNEL_TEX(uint4, __tri_vindex)
|
||||
KERNEL_TEX(uint, __tri_patch)
|
||||
KERNEL_TEX(float2, __tri_patch_uv)
|
||||
KERNEL_TEX(packed_float3, __tri_verts)
|
||||
|
||||
/* curves */
|
||||
KERNEL_TEX(KernelCurve, __curves)
|
||||
KERNEL_TEX(float4, __curve_keys)
|
||||
KERNEL_TEX(KernelCurveSegment, __curve_segments)
|
||||
|
||||
/* patches */
|
||||
KERNEL_TEX(uint, __patches)
|
||||
|
||||
/* pointclouds */
|
||||
KERNEL_TEX(float4, __points)
|
||||
KERNEL_TEX(uint, __points_shader)
|
||||
|
||||
/* attributes */
|
||||
KERNEL_TEX(AttributeMap, __attributes_map)
|
||||
KERNEL_TEX(float, __attributes_float)
|
||||
KERNEL_TEX(float2, __attributes_float2)
|
||||
KERNEL_TEX(packed_float3, __attributes_float3)
|
||||
KERNEL_TEX(float4, __attributes_float4)
|
||||
KERNEL_TEX(uchar4, __attributes_uchar4)
|
||||
|
||||
/* lights */
|
||||
KERNEL_TEX(KernelLightDistribution, __light_distribution)
|
||||
KERNEL_TEX(KernelLight, __lights)
|
||||
KERNEL_TEX(float2, __light_background_marginal_cdf)
|
||||
KERNEL_TEX(float2, __light_background_conditional_cdf)
|
||||
|
||||
/* particles */
|
||||
KERNEL_TEX(KernelParticle, __particles)
|
||||
|
||||
/* shaders */
|
||||
KERNEL_TEX(uint4, __svm_nodes)
|
||||
KERNEL_TEX(KernelShader, __shaders)
|
||||
|
||||
/* lookup tables */
|
||||
KERNEL_TEX(float, __lookup_table)
|
||||
|
||||
/* sobol */
|
||||
KERNEL_TEX(float, __sample_pattern_lut)
|
||||
|
||||
/* image textures */
|
||||
KERNEL_TEX(TextureInfo, __texture_info)
|
||||
|
||||
/* ies lights */
|
||||
KERNEL_TEX(float, __ies)
|
||||
|
||||
#undef KERNEL_TEX
|
@@ -15,11 +15,11 @@ ccl_device float lookup_table_read(KernelGlobals kg, float x, int offset, int si
|
||||
int nindex = min(index + 1, size - 1);
|
||||
float t = x - index;
|
||||
|
||||
float data0 = kernel_tex_fetch(__lookup_table, index + offset);
|
||||
float data0 = kernel_data_fetch(lookup_table, index + offset);
|
||||
if (t == 0.0f)
|
||||
return data0;
|
||||
|
||||
float data1 = kernel_tex_fetch(__lookup_table, nindex + offset);
|
||||
float data1 = kernel_data_fetch(lookup_table, nindex + offset);
|
||||
return (1.0f - t) * data0 + t * data1;
|
||||
}
|
||||
|
||||
|
@@ -152,7 +152,7 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
|
||||
|
||||
KernelFilm *kfilm = &dscene->data.film;
|
||||
|
||||
/* update __data */
|
||||
/* update data */
|
||||
kfilm->exposure = exposure;
|
||||
kfilm->pass_alpha_threshold = pass_alpha_threshold;
|
||||
kfilm->pass_flag = 0;
|
||||
|
@@ -1957,7 +1957,7 @@ void GeometryManager::device_update(Device *device,
|
||||
|
||||
{
|
||||
/* Copy constant data needed by shader evaluation. */
|
||||
device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
|
||||
device->const_copy_to("data", &dscene->data, sizeof(dscene->data));
|
||||
|
||||
scoped_callback_timer timer([scene](double time) {
|
||||
if (scene->update_stats) {
|
||||
|
@@ -697,7 +697,7 @@ void ImageManager::device_load_image(Device *device, Scene *scene, int slot, Pro
|
||||
ImageDataType type = img->metadata.type;
|
||||
|
||||
/* Name for debugging. */
|
||||
img->mem_name = string_printf("__tex_image_%s_%03d", name_from_type(type), slot);
|
||||
img->mem_name = string_printf("tex_image_%s_%03d", name_from_type(type), slot);
|
||||
|
||||
/* Free previous texture in slot. */
|
||||
if (img->mem) {
|
||||
|
@@ -34,7 +34,7 @@ static void shade_background_pixels(Device *device,
|
||||
Progress &progress)
|
||||
{
|
||||
/* Needs to be up to data for attribute access. */
|
||||
device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
|
||||
device->const_copy_to("data", &dscene->data, sizeof(dscene->data));
|
||||
|
||||
const int size = width * height;
|
||||
const int num_channels = 3;
|
||||
|
@@ -34,49 +34,49 @@
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
DeviceScene::DeviceScene(Device *device)
|
||||
: bvh_nodes(device, "__bvh_nodes", MEM_GLOBAL),
|
||||
bvh_leaf_nodes(device, "__bvh_leaf_nodes", MEM_GLOBAL),
|
||||
object_node(device, "__object_node", MEM_GLOBAL),
|
||||
prim_type(device, "__prim_type", MEM_GLOBAL),
|
||||
prim_visibility(device, "__prim_visibility", MEM_GLOBAL),
|
||||
prim_index(device, "__prim_index", MEM_GLOBAL),
|
||||
prim_object(device, "__prim_object", MEM_GLOBAL),
|
||||
prim_time(device, "__prim_time", MEM_GLOBAL),
|
||||
tri_verts(device, "__tri_verts", MEM_GLOBAL),
|
||||
tri_shader(device, "__tri_shader", MEM_GLOBAL),
|
||||
tri_vnormal(device, "__tri_vnormal", MEM_GLOBAL),
|
||||
tri_vindex(device, "__tri_vindex", MEM_GLOBAL),
|
||||
tri_patch(device, "__tri_patch", MEM_GLOBAL),
|
||||
tri_patch_uv(device, "__tri_patch_uv", MEM_GLOBAL),
|
||||
curves(device, "__curves", MEM_GLOBAL),
|
||||
curve_keys(device, "__curve_keys", MEM_GLOBAL),
|
||||
curve_segments(device, "__curve_segments", MEM_GLOBAL),
|
||||
patches(device, "__patches", MEM_GLOBAL),
|
||||
points(device, "__points", MEM_GLOBAL),
|
||||
points_shader(device, "__points_shader", MEM_GLOBAL),
|
||||
objects(device, "__objects", MEM_GLOBAL),
|
||||
object_motion_pass(device, "__object_motion_pass", MEM_GLOBAL),
|
||||
object_motion(device, "__object_motion", MEM_GLOBAL),
|
||||
object_flag(device, "__object_flag", MEM_GLOBAL),
|
||||
object_volume_step(device, "__object_volume_step", MEM_GLOBAL),
|
||||
object_prim_offset(device, "__object_prim_offset", MEM_GLOBAL),
|
||||
camera_motion(device, "__camera_motion", MEM_GLOBAL),
|
||||
attributes_map(device, "__attributes_map", MEM_GLOBAL),
|
||||
attributes_float(device, "__attributes_float", MEM_GLOBAL),
|
||||
attributes_float2(device, "__attributes_float2", MEM_GLOBAL),
|
||||
attributes_float3(device, "__attributes_float3", MEM_GLOBAL),
|
||||
attributes_float4(device, "__attributes_float4", MEM_GLOBAL),
|
||||
attributes_uchar4(device, "__attributes_uchar4", MEM_GLOBAL),
|
||||
light_distribution(device, "__light_distribution", MEM_GLOBAL),
|
||||
lights(device, "__lights", MEM_GLOBAL),
|
||||
light_background_marginal_cdf(device, "__light_background_marginal_cdf", MEM_GLOBAL),
|
||||
light_background_conditional_cdf(device, "__light_background_conditional_cdf", MEM_GLOBAL),
|
||||
particles(device, "__particles", MEM_GLOBAL),
|
||||
svm_nodes(device, "__svm_nodes", MEM_GLOBAL),
|
||||
shaders(device, "__shaders", MEM_GLOBAL),
|
||||
lookup_table(device, "__lookup_table", MEM_GLOBAL),
|
||||
sample_pattern_lut(device, "__sample_pattern_lut", MEM_GLOBAL),
|
||||
ies_lights(device, "__ies", MEM_GLOBAL)
|
||||
: bvh_nodes(device, "bvh_nodes", MEM_GLOBAL),
|
||||
bvh_leaf_nodes(device, "bvh_leaf_nodes", MEM_GLOBAL),
|
||||
object_node(device, "object_node", MEM_GLOBAL),
|
||||
prim_type(device, "prim_type", MEM_GLOBAL),
|
||||
prim_visibility(device, "prim_visibility", MEM_GLOBAL),
|
||||
prim_index(device, "prim_index", MEM_GLOBAL),
|
||||
prim_object(device, "prim_object", MEM_GLOBAL),
|
||||
prim_time(device, "prim_time", MEM_GLOBAL),
|
||||
tri_verts(device, "tri_verts", MEM_GLOBAL),
|
||||
tri_shader(device, "tri_shader", MEM_GLOBAL),
|
||||
tri_vnormal(device, "tri_vnormal", MEM_GLOBAL),
|
||||
tri_vindex(device, "tri_vindex", MEM_GLOBAL),
|
||||
tri_patch(device, "tri_patch", MEM_GLOBAL),
|
||||
tri_patch_uv(device, "tri_patch_uv", MEM_GLOBAL),
|
||||
curves(device, "curves", MEM_GLOBAL),
|
||||
curve_keys(device, "curve_keys", MEM_GLOBAL),
|
||||
curve_segments(device, "curve_segments", MEM_GLOBAL),
|
||||
patches(device, "patches", MEM_GLOBAL),
|
||||
points(device, "points", MEM_GLOBAL),
|
||||
points_shader(device, "points_shader", MEM_GLOBAL),
|
||||
objects(device, "objects", MEM_GLOBAL),
|
||||
object_motion_pass(device, "object_motion_pass", MEM_GLOBAL),
|
||||
object_motion(device, "object_motion", MEM_GLOBAL),
|
||||
object_flag(device, "object_flag", MEM_GLOBAL),
|
||||
object_volume_step(device, "object_volume_step", MEM_GLOBAL),
|
||||
object_prim_offset(device, "object_prim_offset", MEM_GLOBAL),
|
||||
camera_motion(device, "camera_motion", MEM_GLOBAL),
|
||||
attributes_map(device, "attributes_map", MEM_GLOBAL),
|
||||
attributes_float(device, "attributes_float", MEM_GLOBAL),
|
||||
attributes_float2(device, "attributes_float2", MEM_GLOBAL),
|
||||
attributes_float3(device, "attributes_float3", MEM_GLOBAL),
|
||||
attributes_float4(device, "attributes_float4", MEM_GLOBAL),
|
||||
attributes_uchar4(device, "attributes_uchar4", MEM_GLOBAL),
|
||||
light_distribution(device, "light_distribution", MEM_GLOBAL),
|
||||
lights(device, "lights", MEM_GLOBAL),
|
||||
light_background_marginal_cdf(device, "light_background_marginal_cdf", MEM_GLOBAL),
|
||||
light_background_conditional_cdf(device, "light_background_conditional_cdf", MEM_GLOBAL),
|
||||
particles(device, "particles", MEM_GLOBAL),
|
||||
svm_nodes(device, "svm_nodes", MEM_GLOBAL),
|
||||
shaders(device, "shaders", MEM_GLOBAL),
|
||||
lookup_table(device, "lookup_table", MEM_GLOBAL),
|
||||
sample_pattern_lut(device, "sample_pattern_lut", MEM_GLOBAL),
|
||||
ies_lights(device, "ies", MEM_GLOBAL)
|
||||
{
|
||||
memset((void *)&data, 0, sizeof(data));
|
||||
}
|
||||
@@ -366,7 +366,7 @@ void Scene::device_update(Device *device_, Progress &progress)
|
||||
dscene.data.volume_stack_size = get_volume_stack_size();
|
||||
|
||||
progress.set_status("Updating Device", "Writing constant memory");
|
||||
device->const_copy_to("__data", &dscene.data, sizeof(dscene.data));
|
||||
device->const_copy_to("data", &dscene.data, sizeof(dscene.data));
|
||||
}
|
||||
|
||||
if (print_stats) {
|
||||
|
Reference in New Issue
Block a user