[Bf-blender-cvs] [4ca688a] master: Cycles: OpenCL split kernel cleanup, move casts from .h files to .cl files
Sergey Sharybin
noreply at git.blender.org
Thu Oct 29 17:53:00 CET 2015
Commit: 4ca688a963408fe326ba8e2a696d7b0fdc4eb1e4
Author: Sergey Sharybin
Date: Thu Oct 29 21:44:36 2015 +0500
Branches: master
https://developer.blender.org/rB4ca688a963408fe326ba8e2a696d7b0fdc4eb1e4
Cycles: OpenCL split kernel cleanup, move casts from .h files to .cl files
Ideally we shouldn't use char* at all, but for now we have to, so at least
let's assume common .h files are free from pointer magic.
===================================================================
M intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
M intern/cycles/kernel/kernels/opencl/kernel_data_init.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/split/kernel_background_buffer_update.h
M intern/cycles/kernel/split/kernel_data_init.h
M intern/cycles/kernel/split/kernel_direct_lighting.h
M intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
M intern/cycles/kernel/split/kernel_lamp_emission.h
M intern/cycles/kernel/split/kernel_next_iteration_setup.h
M intern/cycles/kernel/split/kernel_scene_intersect.h
M intern/cycles/kernel/split/kernel_shader_eval.h
M intern/cycles/kernel/split/kernel_shadow_blocked.h
===================================================================
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 eff77b8..c5eb7a2 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
@@ -17,9 +17,9 @@
#include "split/kernel_background_buffer_update.h"
__kernel void kernel_ocl_path_trace_background_buffer_update(
- ccl_global char *globals,
+ ccl_global char *kg,
ccl_constant KernelData *data,
- ccl_global char *shader_data,
+ ccl_global char *sd,
ccl_global float *per_sample_output_buffers,
ccl_global uint *rng_state,
ccl_global uint *rng_coop, /* Required for buffer Update */
@@ -83,9 +83,9 @@ __kernel void kernel_ocl_path_trace_background_buffer_update(
if(ray_index != QUEUE_EMPTY_SLOT) {
#endif
enqueue_flag =
- kernel_background_buffer_update(globals,
+ kernel_background_buffer_update((KernelGlobals *)kg,
data,
- shader_data,
+ (ShaderData *)sd,
per_sample_output_buffers,
rng_state,
rng_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
index c327767..8329aa9 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -17,9 +17,9 @@
#include "split/kernel_data_init.h"
__kernel void kernel_ocl_path_trace_data_init(
- ccl_global char *globals,
- ccl_global char *shader_data_sd, /* Arguments related to ShaderData */
- ccl_global char *shader_data_sd_DL_shadow, /* Arguments related to ShaderData */
+ ccl_global char *kg,
+ ccl_global char *sd,
+ ccl_global char *sd_DL_shadow,
ccl_global float3 *P_sd,
ccl_global float3 *P_sd_DL_shadow,
@@ -141,9 +141,9 @@ __kernel void kernel_ocl_path_trace_data_init(
#endif
int parallel_samples) /* Number of samples to be processed in parallel */
{
- kernel_data_init(globals,
- shader_data_sd,
- shader_data_sd_DL_shadow,
+ kernel_data_init((KernelGlobals *)kg,
+ (ShaderData *)sd,
+ (ShaderData *)sd_DL_shadow,
P_sd,
P_sd_DL_shadow,
N_sd,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
index 6ec7501..83e0afd 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
@@ -17,10 +17,10 @@
#include "split/kernel_direct_lighting.h"
__kernel void kernel_ocl_path_trace_direct_lighting(
- ccl_global char *globals,
+ ccl_global char *kg,
ccl_constant KernelData *data,
- ccl_global char *shader_data, /* Required for direct lighting */
- ccl_global char *shader_DL, /* Required for direct lighting */
+ 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 */
ccl_global PathState *PathState_coop, /* Required for direct lighting */
ccl_global int *ISLamp_coop, /* Required for direct lighting */
@@ -61,10 +61,10 @@ __kernel void kernel_ocl_path_trace_direct_lighting(
#ifndef __COMPUTE_DEVICE_GPU__
if(ray_index != QUEUE_EMPTY_SLOT) {
#endif
- enqueue_flag = kernel_direct_lighting(globals,
+ enqueue_flag = kernel_direct_lighting((KernelGLobals *)kg,
data,
- shader_data,
- shader_DL,
+ (ShaderData *)sd,
+ (ShaderData *)sd_DL,
rng_coop,
PathState_coop,
ISLamp_coop,
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 ae5f5cd..6fe25d7 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
@@ -17,9 +17,9 @@
#include "split/kernel_holdout_emission_blurring_pathtermination_ao.h"
__kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
- ccl_global char *globals,
+ ccl_global char *kg,
ccl_constant KernelData *data,
- ccl_global char *shader_data, /* Required throughout the kernel except probabilistic path termination and AO */
+ 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 */
ccl_global float3 *throughput_coop, /* Required for handling holdout material and AO */
@@ -75,9 +75,9 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao
if(ray_index != QUEUE_EMPTY_SLOT) {
#endif
kernel_holdout_emission_blurring_pathtermination_ao(
- globals,
+ (KernelGlobals *)kg,
data,
- shader_data,
+ (ShaderData *)sd,
per_sample_output_buffers,
rng_coop,
throughput_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
index 1bc7808..fb327a3 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
@@ -17,9 +17,9 @@
#include "split/kernel_lamp_emission.h"
__kernel void kernel_ocl_path_trace_lamp_emission(
- ccl_global char *globals,
+ ccl_global char *kg,
ccl_constant KernelData *data,
- ccl_global char *shader_data, /* Required for lamp emission */
+ ccl_global char *sd, /* Required for lamp emission */
ccl_global float3 *throughput_coop, /* Required for lamp emission */
PathRadiance *PathRadiance_coop, /* Required for lamp emission */
ccl_global Ray *Ray_coop, /* Required for lamp emission */
@@ -68,9 +68,9 @@ __kernel void kernel_ocl_path_trace_lamp_emission(
}
}
- kernel_lamp_emission(globals,
+ kernel_lamp_emission((KenrelGLobals *)kg,
data,
- shader_data,
+ (ShaderData *)sd,
throughput_coop,
PathRadiance_coop,
Ray_coop,
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 dcf4db4..8896fe7 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
@@ -17,9 +17,9 @@
#include "split/kernel_next_iteration_setup.h"
__kernel void kernel_ocl_path_trace_next_iteration_setup(
- ccl_global char *globals,
+ ccl_global char *kg,
ccl_constant KernelData *data,
- ccl_global char *shader_data, /* Required for setting up ray for next iteration */
+ 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 */
PathRadiance *PathRadiance_coop, /* Required for setting up ray for next iteration */
@@ -83,9 +83,9 @@ __kernel void kernel_ocl_path_trace_next_iteration_setup(
#ifndef __COMPUTE_DEVICE_GPU__
if(ray_index != QUEUE_EMPTY_SLOT) {
#endif
- enqueue_flag = kernel_next_iteration_setup(globals,
+ enqueue_flag = kernel_next_iteration_setup((KernelGlobals *)kg,
data,
- shader_data,
+ (ShaderData *)sd,
rng_coop,
throughput_coop,
PathRadiance_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
index e5fad7b..3750a0d 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
@@ -17,7 +17,7 @@
#include "split/kernel_scene_intersect.h"
__kernel void kernel_ocl_path_trace_scene_intersect(
- ccl_global char *globals,
+ ccl_global char *kg,
ccl_constant KernelData *data,
ccl_global uint *rng_coop,
ccl_global Ray *Ray_coop, /* Required for scene_intersect */
@@ -65,7 +65,7 @@ __kernel void kernel_ocl_path_trace_scene_intersect(
}
}
- kernel_scene_intersect(globals,
+ kernel_scene_intersect((KernelGlobals *)kg,
data,
rng_coop,
Ray_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
index b9f616e..2fa5496 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
+++ b/inte
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list