[Bf-blender-cvs] [9bce104] master: Cycles: Partially revert previous commit

Sergey Sharybin noreply at git.blender.org
Sun Nov 1 17:05:56 CET 2015


Commit: 9bce104c8c72cdd4deac8e59ce571892302fb366
Author: Sergey Sharybin
Date:   Sun Nov 1 21:00:26 2015 +0500
Branches: master
https://developer.blender.org/rB9bce104c8c72cdd4deac8e59ce571892302fb366

Cycles: Partially revert previous commit

Apparently removing kernel arguments broke NVidia OpenCL.

Needs more investigation, for the time being revering changes which caused problem.

===================================================================

M	intern/cycles/device/device_opencl.cpp
M	intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
M	intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
M	intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
M	intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
M	intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
M	intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
M	intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
M	intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
M	intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl
M	intern/cycles/kernel/split/kernel_sum_all_radiance.h

===================================================================

diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index b4fc41a..d6b1ab3 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -2843,6 +2843,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_scene_intersect,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                rng_coop,
 		                Ray_coop,
 		                PathState_coop,
@@ -2862,6 +2863,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_lamp_emission,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                sd,
 		                throughput_coop,
 		                PathRadiance_coop,
@@ -2887,6 +2889,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_background_buffer_update,
 		                 0,
 		                 kgbuffer,
+		                 d_data,
 		                 sd,
 		                 per_sample_output_buffers,
 		                 d_rng_state,
@@ -2923,6 +2926,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_shader_eval,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                sd,
 		                rng_coop,
 		                Ray_coop,
@@ -2936,6 +2940,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                sd,
 		                per_sample_output_buffers,
 		                rng_coop,
@@ -2965,6 +2970,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_direct_lighting,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                sd,
 		                sd_DL_shadow,
 		                rng_coop,
@@ -2980,6 +2986,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_shadow_blocked,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                sd_DL_shadow,
 		                PathState_coop,
 		                LightRay_coop,
@@ -2995,6 +3002,7 @@ public:
 		kernel_set_args(ckPathTraceKernel_next_iteration_setup,
 		                0,
 		                kgbuffer,
+		                d_data,
 		                sd,
 		                rng_coop,
 		                throughput_coop,
@@ -3015,7 +3023,7 @@ public:
 
 		kernel_set_args(ckPathTraceKernel_sum_all_radiance,
 		                0,
-		                kgbuffer,
+		                d_data,
 		                d_buffer,
 		                per_sample_output_buffers,
 		                num_parallel_samples,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
index c28338b..a3eecd3 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_background_buffer_update(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd,
         ccl_global float *per_sample_output_buffers,
         ccl_global uint *rng_state,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
index a646d42..d4a7cff 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_direct_lighting(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd,                    /* Required for direct lighting */
         ccl_global char *sd_DL,                 /* Required for direct lighting */
         ccl_global uint *rng_coop,              /* Required for direct lighting */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
index 2f6d3e2..e063614 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd,                   /* Required throughout the kernel except probabilistic path termination and AO */
         ccl_global float *per_sample_output_buffers,
         ccl_global uint *rng_coop,             /* Required for "kernel_write_data_passes" and AO */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
index 73de387..5215a0e 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_lamp_emission(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd,                   /* Required for lamp emission */
         ccl_global float3 *throughput_coop,    /* Required for lamp emission */
         PathRadiance *PathRadiance_coop,       /* Required for lamp emission */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
index f6962aa..6d49b62 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_next_iteration_setup(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd,                  /* Required for setting up ray for next iteration */
         ccl_global uint *rng_coop,            /* Required for setting up ray for next iteration */
         ccl_global float3 *throughput_coop,   /* Required for setting up ray for next iteration */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
index c117008..12eff6c 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_scene_intersect(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global uint *rng_coop,
         ccl_global Ray *Ray_coop,              /* Required for scene_intersect */
         ccl_global PathState *PathState_coop,  /* Required for scene_intersect */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
index 6f9faf5..c37856c 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_shader_eval(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd,                   /* Output ShaderData structure to be filled */
         ccl_global uint *rng_coop,             /* Required for rbsdf calculation */
         ccl_global Ray *Ray_coop,              /* Required for setting up shader from ray */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
index b2ee6d2..260a601 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
@@ -18,6 +18,7 @@
 
 __kernel void kernel_ocl_path_trace_shadow_blocked(
         ccl_global char *kg,
+        ccl_constant KernelData *data,
         ccl_global char *sd_shadow,            /* Required for shadow blocked */
         ccl_global PathState *PathState_coop,  /* Required for shadow blocked */
         ccl_global Ray *LightRay_dl_coop,      /* Required for direct lighting's shadow blocked */
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl b/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl
index 0ebfae9..88a1ed8 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl
@@ -17,7 +17,7 @@
 #include "split/kernel_sum_all_radiance.h"
 
 __kernel void kernel_ocl_path_trace_sum_all_radiance(
-        ccl_global char *kg,                         /* To get pass_stride to offet into buffer */
+        ccl_constant KernelData *data,               /* To get pass_stride to offet into buffer */
         ccl_global float *buffer,                    /* Output buffer of RenderTile */
         ccl_global float *per_sample_output_buffer,  /* Radiance contributed by all samples */
         int parallel_samples, int sw, int sh, int stride,
@@ -26,7 +26,7 @@ __kernel void kernel_ocl_path_trace_sum_all_radiance(
         int buffer_stride,
         int start_sample)
 {
-	kernel_sum_all_radiance((KernelGlobals *)kg,
+	kernel_sum_all_radiance(data,
 	                        buffer,
 	                        per_sample_output_buffer,
 	                        parallel_samples,
diff --git a/intern/cycles/kernel/split/kernel_sum_all_radiance.h b/intern/cycles/kernel/split/kernel_sum_all_radiance.h
index c860fc4..a21e9b6 100644
--- a/intern/cycles/kernel/split/kernel_sum_all_radiance.h
+++ b/intern/cycles/kernel/split/kernel_sum_all_radiance.h
@@ -24,7 +24,7 @@
  * by all different samples and stores them in the RenderTile's output buffer.
  */
 ccl_device void kernel_sum_all_radiance(
-        KernelGlobals *kg,                           /* To get pass_stride to offet into buffer */
+        ccl_constant KernelData *data,               /* To get pass_stride to offet into buffer */
         ccl_global float *buffer,                    /* Output buffer of RenderTile */
         ccl_global float *per_sample_

@@ Diff output truncated at 10240 characters. @@




More information about the Bf-blender-cvs mailing list