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. @@
_______________________________________________
Bf-blender-cvs mailing list
[email protected]
http://lists.blender.org/mailman/listinfo/bf-blender-cvs
