From 9bce104c8c72cdd4deac8e59ce571892302fb366 Mon Sep 17 00:00:00 2001 From: Sergey Sharybin Date: Sun, 1 Nov 2015 21:00:26 +0500 Subject: Cycles: Partially revert previous commit Apparently removing kernel arguments broke NVidia OpenCL. Needs more investigation, for the time being revering changes which caused problem. --- intern/cycles/device/device_opencl.cpp | 10 +++++++++- .../kernel/kernels/opencl/kernel_background_buffer_update.cl | 1 + intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl | 1 + .../kernel_holdout_emission_blurring_pathtermination_ao.cl | 1 + intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl | 1 + .../kernel/kernels/opencl/kernel_next_iteration_setup.cl | 1 + intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl | 1 + intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl | 1 + intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl | 1 + intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl | 4 ++-- intern/cycles/kernel/split/kernel_sum_all_radiance.h | 10 +++++----- 11 files changed, 24 insertions(+), 8 deletions(-) diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp index b4fc41ab247..d6b1ab38bd0 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 c28338be5f2..a3eecd3128b 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 a646d42193d..d4a7cffb403 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 2f6d3e2d404..e063614da1a 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 73de387ba86..5215a0e0827 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 f6962aad9e8..6d49b6294a8 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 c117008b99c..12eff6ccc1a 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 6f9faf509ac..c37856c8f30 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 b2ee6d2b776..260a601946f 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 0ebfae911a8..88a1ed830af 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 c860fc42a86..a21e9b6a0b1 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_output_buffer, /* Radiance contributed by all samples */ int parallel_samples, int sw, int sh, int stride, @@ -37,14 +37,14 @@ ccl_device void kernel_sum_all_radiance( int y = get_global_id(1); if(x < sw && y < sh) { - buffer += ((buffer_offset_x + x) + (buffer_offset_y + y) * buffer_stride) * (kg->data->film.pass_stride); - per_sample_output_buffer += ((x + y * stride) * parallel_samples) * (kg->data->film.pass_stride); + buffer += ((buffer_offset_x + x) + (buffer_offset_y + y) * buffer_stride) * (data->film.pass_stride); + per_sample_output_buffer += ((x + y * stride) * parallel_samples) * (data->film.pass_stride); - int sample_stride = (kg->data->film.pass_stride); + int sample_stride = (data->film.pass_stride); int sample_iterator = 0; int pass_stride_iterator = 0; - int num_floats = kg->data->film.pass_stride; + int num_floats = data->film.pass_stride; for(sample_iterator = 0; sample_iterator < parallel_samples; sample_iterator++) { for(pass_stride_iterator = 0; pass_stride_iterator < num_floats; pass_stride_iterator++) { -- cgit v1.2.3