[Bf-blender-cvs] [e93ee69] cycles_kernel_split: [BCYCLES-209] make direct lighting and shadow blocked kernels share shaderData global buffers
Kavitha Sampath
noreply at git.blender.org
Wed Apr 15 17:37:10 CEST 2015
Commit: e93ee697a94bc5880b351ed5b1ca97fd60a3165c
Author: Kavitha Sampath
Date: Fri Apr 10 18:30:45 2015 +0530
Branches: cycles_kernel_split
https://developer.blender.org/rBe93ee697a94bc5880b351ed5b1ca97fd60a3165c
[BCYCLES-209] make direct lighting and shadow blocked kernels share shaderData global buffers
===================================================================
M intern/cycles/device/device_opencl.cpp
M intern/cycles/kernel/kernel_DataInit.cl
===================================================================
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index d9bfdd2..b6fd5a0 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -411,8 +411,7 @@ public:
/* global buffers for ShaderData */
cl_mem sd; /* ShaderData used in the main path-iteration loop */
- cl_mem sd_dl; /* ShaderData used in DirectLighting kernel */
- cl_mem sd_shadow; /* ShaderData used in ShadowBlocked kernel */
+ cl_mem sd_DL_shadow; /* ShaderData used in Direct Lighting and ShadowBlocked kernel */
/* global buffers of each member of ShaderData */
cl_mem P_sd;
@@ -674,8 +673,7 @@ public:
/* Initialize cl_mem variables */
kgbuffer = NULL;
sd = NULL;
- sd_dl = NULL;
- sd_shadow = NULL;
+ sd_DL_shadow = NULL;
P_sd = NULL;
P_sd_dl = NULL;;
@@ -2164,11 +2162,8 @@ public:
if(sd != NULL)
clReleaseMemObject(sd);
- if(sd_dl != NULL)
- clReleaseMemObject(sd_dl);
-
- if(sd_shadow != NULL)
- clReleaseMemObject(sd_shadow);
+ if(sd_DL_shadow != NULL)
+ clReleaseMemObject(sd_DL_shadow);
if(ray_state != NULL)
clReleaseMemObject(ray_state);
@@ -2499,11 +2494,8 @@ public:
sd = clCreateBuffer(cxContext, CL_MEM_READ_WRITE, get_shaderdata_soa_size(), NULL, &ciErr);
assert(ciErr == CL_SUCCESS && "Can't create Shaderdata memory");
- sd_dl = clCreateBuffer(cxContext, CL_MEM_READ_WRITE, get_shaderdata_soa_size(), NULL, &ciErr);
- assert(ciErr == CL_SUCCESS && "Can't create sd_dl memory");
-
- sd_shadow = clCreateBuffer(cxContext, CL_MEM_READ_WRITE, get_shaderdata_soa_size(), NULL, &ciErr);
- assert(ciErr == CL_SUCCESS && "Can't create sd_shadow memory");
+ sd_DL_shadow = clCreateBuffer(cxContext, CL_MEM_READ_WRITE, get_shaderdata_soa_size(), NULL, &ciErr);
+ assert(ciErr == CL_SUCCESS && "Can't create sd_DL_shadow memory");
P_sd = clCreateBuffer(cxContext, CL_MEM_READ_WRITE, num_global_elements * sizeof(float3), NULL, &ciErr);
assert(ciErr == CL_SUCCESS && "Can't create P_sd memory");
@@ -2811,8 +2803,7 @@ public:
opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(kgbuffer), (void*)&kgbuffer));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(sd), (void*)&sd));
- opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(sd_dl), (void*)&sd_dl));
- opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(sd_shadow), (void*)&sd_shadow));
+ opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(sd_DL_shadow), (void*)&sd_DL_shadow));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(P_sd), (void*)&P_sd));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(P_sd_dl), (void*)&P_sd_dl));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DataInit_SPLIT_KERNEL, narg++, sizeof(P_sd_shadow), (void*)&P_sd_shadow));
@@ -3103,7 +3094,7 @@ public:
opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(kgbuffer), (void*)&kgbuffer));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(d_data), (void*)&d_data));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(sd), (void*)&sd));
- opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(sd_dl), (void*)&sd_dl));
+ opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(sd_DL_shadow), (void*)&sd_DL_shadow));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(rng_coop), (void*)&rng_coop));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(PathState_coop), (void*)&PathState_coop));
opencl_assert(clSetKernelArg(ckPathTraceKernel_DirectLighting_SPLIT_KERNEL, narg++, sizeof(ISLamp_coop), (void*)&ISLamp_coop));
@@ -3119,7 +3110,7 @@ public:
narg = 0;
opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(kgbuffer), (void*)&kgbuffer));
opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(d_data), (void*)&d_data));
- opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(sd_shadow), (void*)&sd_shadow));
+ opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(sd_DL_shadow), (void*)&sd_DL_shadow));
opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(PathState_coop), (void*)&PathState_coop));
opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(LightRay_coop), (void*)&LightRay_coop));
opencl_assert(clSetKernelArg(ckPathTraceKernel_ShadowBlocked_DirectLighting_SPLIT_KERNEL, narg++, sizeof(AOLightRay_coop), (void*)&AOLightRay_coop));
@@ -3394,8 +3385,7 @@ public:
total_invariable_mem_allocated += NUM_QUEUES * sizeof(unsigned int); /* Queue index size */
total_invariable_mem_allocated += sizeof(char); /* use_queues_flag size */
total_invariable_mem_allocated += ShaderData_SOA_size; /* sd size */
- total_invariable_mem_allocated += ShaderData_SOA_size; /* sd_dl size */
- total_invariable_mem_allocated += ShaderData_SOA_size; /* sd_shadow size */
+ total_invariable_mem_allocated += ShaderData_SOA_size; /* sd_DL_shadow size */
return total_invariable_mem_allocated;
}
@@ -3458,7 +3448,6 @@ public:
+ Intersection_coop_AO_size /* Instersection_coop_AO */
+ Intersection_coop_DL_size /* Intersection coop DL */
+ shaderdata_volume /* Overall ShaderData */
- + shaderdata_volume /* ShaderData_coop_DL */
+ (shaderdata_volume * 2) /* ShaderData coop shadow */
+ LightRay_size + BSDFEval_size + AOAlpha_size + AOBSDF_size + AOLightRay_size
+ (sizeof(int)* NUM_QUEUES)
diff --git a/intern/cycles/kernel/kernel_DataInit.cl b/intern/cycles/kernel/kernel_DataInit.cl
index 7964f9c..a6fbaf3 100644
--- a/intern/cycles/kernel/kernel_DataInit.cl
+++ b/intern/cycles/kernel/kernel_DataInit.cl
@@ -80,7 +80,6 @@ ccl_device_inline void kernel_path_trace_setup_SPLIT_KERNEL(__ADDR_SPACE__ Kerne
__kernel void kernel_ocl_path_trace_data_initialization_SPLIT_KERNEL(
ccl_global char *globals,
ccl_global char *shader_data_sd, /* Arguments related to ShaderData */
- ccl_global char *shader_data_sd_dl, /* Arguments related to ShaderData */
ccl_global char *shader_data_sd_shadow, /* Arguments related to ShaderData */
ccl_global float3 *P_sd,
@@ -230,114 +229,87 @@ __kernel void kernel_ocl_path_trace_data_initialization_SPLIT_KERNEL(
/* Load ShaderData structure */
ccl_global ShaderData *sd = (ccl_global ShaderData *)shader_data_sd;
- ccl_global ShaderData *sd_dl = (ccl_global ShaderData *)shader_data_sd_dl;
ccl_global ShaderData *sd_shadow = (ccl_global ShaderData *)shader_data_sd_shadow;
sd->P = P_sd;
- sd_dl->P = P_sd_dl;
sd_shadow->P = P_sd_shadow;
sd->N = N_sd;
- sd_dl->N = N_sd_dl;
sd_shadow->N = N_sd_shadow;
sd->Ng = Ng_sd;
- sd_dl->Ng = Ng_sd_dl;
sd_shadow->Ng = Ng_sd_shadow;
sd->I = I_sd;
- sd_dl->I = I_sd_dl;
sd_shadow->I = I_sd_shadow;
sd->shader = shader_sd;
- sd_dl->shader = shader_sd_dl;
sd_shadow->shader = shader_sd_shadow;
sd->flag = flag_sd;
- sd_dl->flag = flag_sd_dl;
sd_shadow->flag = flag_sd_shadow;
sd->prim = prim_sd;
- sd_dl->prim = prim_sd_dl;
sd_shadow->prim = prim_sd_shadow;
sd->type = type_sd;
- sd_dl->type = type_sd_dl;
sd_shadow->type = type_sd_shadow;
sd->u = u_sd;
- sd_dl->u = u_sd_dl;
sd_shadow->u = u_sd_shadow;
sd->v = v_sd;
- sd_dl->v = v_sd_dl;
sd_shadow->v = v_sd_shadow;
sd->object = object_sd;
- sd_dl->object = object_sd_dl;
sd_shadow->object = object_sd_shadow;
sd->time = time_sd;
- sd_dl->time = time_sd_dl;
sd_shadow->time = time_sd_shadow;
sd->ray_length = ray_length_sd;
- sd_dl->ray_length = ray_length_sd_dl;
sd_shadow->ray_length = ray_length_sd_shadow;
sd->ray_depth = ray_depth_sd;
- sd_dl->ray_depth = ray_depth_sd_dl;
sd_shadow->ray_depth = ray_depth_sd_shadow;
sd->transparent_depth = transparent_depth_sd;
- sd_dl->transparent_depth = transparent_depth_sd_dl;
sd_shadow->transparent_depth = transparent_depth_sd_shadow;
#ifdef __RAY_DIFFERENTIALS__
sd->dP = dP_sd;
- sd_dl->dP = dP_sd_dl;
sd_shadow->dP = dP_sd_shadow;
sd->dI = dI_sd;
- sd_dl->dI = dI_sd_dl;
sd_shadow->dI = dI_sd_shadow;
sd->du = du_sd;
- sd_dl->du = du_sd_dl;
sd_shadow->du = du_sd_shadow;
sd->dv = dv_sd;
- sd_dl->dv = dv_sd_dl;
sd_shadow->dv = dv_sd_shadow;
#ifdef __DPDU__
sd->dPdu = dPdu_sd;
- sd_dl->dPdu = dPdu_sd_dl;
sd_shadow->dPdu = dPdu_sd_shadow;
sd->dPdv = dPdv_sd;
- sd_dl->dPdv = dPdv_sd_dl;
sd_shadow->dPdv = dPdv_sd_shadow;
#endif
#endif
sd->closure = closure_sd;
- sd_dl->closure = closure_sd_dl;
sd_shadow->closure = closure_sd_shadow;
sd->num_closure = num_closure_sd;
- sd_dl->num_closure = num_closure_sd_dl;
sd_shadow->num_closure = num_closure_sd_shadow;
sd->randb_closure = randb_closure_sd;
- sd_dl->randb_closure = randb_closure_sd_dl;
sd_shadow->randb_closure = randb_closure_sd_shadow;
sd->ray_P = ray_P_sd;
- sd_dl->ray_P = ray_P_sd_dl;
sd_shadow->ray_P = ray_P_sd_shadow;
sd->ray_dP = ray_dP_sd;
- sd_dl->ray_dP = ray_dP_sd_dl;
sd_shadow->ray_dP = ray_dP_sd_shadow;
int thread_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
More information about the Bf-blender-cvs
mailing list