[Bf-blender-cvs] [7e7cbe47f8b] experimental-build: Revert "Cycles: Pass all buffers to each kernel call for OpenCL"
Mai Lavelle
noreply at git.blender.org
Mon Jun 5 17:03:55 CEST 2017
Commit: 7e7cbe47f8b1d834263309d3c24efbb4fb85309c
Author: Mai Lavelle
Date: Mon Jun 5 11:03:06 2017 -0400
Branches: experimental-build
https://developer.blender.org/rB7e7cbe47f8b1d834263309d3c24efbb4fb85309c
Revert "Cycles: Pass all buffers to each kernel call for OpenCL"
This reverts commit 18cc5afa58fc8df01cd3f5353b875374053045ed.
===================================================================
M intern/cycles/device/opencl/opencl_split.cpp
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
M intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
M intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
M intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
M intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
M intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.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_path_init.cl
M intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.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_shader_setup.cl
M intern/cycles/kernel/kernels/opencl/kernel_shader_sort.cl
M intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
M intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
D intern/cycles/kernel/kernels/opencl/kernel_split_function.h
M intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl
===================================================================
diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp
index fbd54918879..76dcbd6fc9a 100644
--- a/intern/cycles/device/opencl/opencl_split.cpp
+++ b/intern/cycles/device/opencl/opencl_split.cpp
@@ -176,62 +176,17 @@ protected:
friend class OpenCLSplitKernelFunction;
};
-struct CachedSplitMemory {
- int id;
- device_memory* split_data;
- device_memory* ray_state;
- device_ptr* rng_state;
- device_memory* queue_index;
- device_memory* use_queues_flag;
- device_memory* work_pools;
- device_ptr* buffer;
-};
-
class OpenCLSplitKernelFunction : public SplitKernelFunction {
public:
OpenCLDeviceSplitKernel* device;
OpenCLDeviceBase::OpenCLProgram program;
- CachedSplitMemory& cached_memory;
- int cached_id;
-
- OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device, CachedSplitMemory& cached_memory) :
- device(device), cached_memory(cached_memory), cached_id(cached_memory.id-1)
- {
- }
- ~OpenCLSplitKernelFunction()
- {
- program.release();
- }
+ OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device) : device(device) {}
+ ~OpenCLSplitKernelFunction() { program.release(); }
virtual bool enqueue(const KernelDimensions& dim, device_memory& kg, device_memory& data)
{
- if(cached_id != cached_memory.id) {
- cl_uint start_arg_index =
- device->kernel_set_args(program(),
- 0,
- kg,
- data,
- *cached_memory.split_data,
- *cached_memory.ray_state,
- *cached_memory.rng_state);
-
-/* TODO(sergey): Avoid map lookup here. */
-#define KERNEL_TEX(type, ttype, name) \
- device->set_kernel_arg_mem(program(), &start_arg_index, #name);
-#include "kernel/kernel_textures.h"
-#undef KERNEL_TEX
-
- start_arg_index +=
- device->kernel_set_args(program(),
- start_arg_index,
- *cached_memory.queue_index,
- *cached_memory.use_queues_flag,
- *cached_memory.work_pools,
- *cached_memory.buffer);
-
- cached_id = cached_memory.id;
- }
+ device->kernel_set_args(program(), 0, kg, data);
device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue,
program(),
@@ -258,7 +213,6 @@ public:
class OpenCLSplitKernel : public DeviceSplitKernel {
OpenCLDeviceSplitKernel *device;
- CachedSplitMemory cached_memory;
public:
explicit OpenCLSplitKernel(OpenCLDeviceSplitKernel *device) : DeviceSplitKernel(device), device(device) {
}
@@ -266,7 +220,7 @@ public:
virtual SplitKernelFunction* get_split_kernel_function(string kernel_name,
const DeviceRequestedFeatures& requested_features)
{
- OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device, cached_memory);
+ OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device);
bool single_program = OpenCLInfo::use_single_program();
kernel->program =
@@ -395,15 +349,6 @@ public:
return false;
}
- cached_memory.split_data = &split_data;
- cached_memory.ray_state = &ray_state;
- cached_memory.rng_state = &rtile.rng_state;
- cached_memory.queue_index = &queue_index;
- cached_memory.use_queues_flag = &use_queues_flag;
- cached_memory.work_pools = &work_pool_wgs;
- cached_memory.buffer = &rtile.buffer;
- cached_memory.id++;
-
return true;
}
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 2d5ee11367d..bef869f34b4 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -121,10 +121,6 @@ set(SRC_KERNELS_CUDA_HEADERS
kernels/cuda/kernel_config.h
)
-set(SRC_KERNELS_OPENCL_HEADERS
- kernels/opencl/kernel_split_function.h
-)
-
set(SRC_CLOSURE_HEADERS
closure/alloc.h
closure/bsdf.h
@@ -454,7 +450,6 @@ add_library(cycles_kernel
${SRC_HEADERS}
${SRC_KERNELS_CPU_HEADERS}
${SRC_KERNELS_CUDA_HEADERS}
- ${SRC_KERNELS_OPENCL_HEADERS}
${SRC_BVH_HEADERS}
${SRC_CLOSURE_HEADERS}
${SRC_FILTER_HEADERS}
@@ -498,7 +493,6 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_shadow_blocke
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_next_iteration_setup.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_indirect_subsurface.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_buffer_update.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
-delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/kernel_split_function.h" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/filter.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/kernel_split.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda)
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
index dcea2630aef..db65c91baf7 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
@@ -18,9 +18,10 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_buffer_update.h"
-#define KERNEL_NAME buffer_update
-#define LOCALS_TYPE unsigned int
-#include "kernel/kernels/opencl/kernel_split_function.h"
-#undef KERNEL_NAME
-#undef LOCALS_TYPE
-
+__kernel void kernel_ocl_path_trace_buffer_update(
+ ccl_global char *kg,
+ ccl_constant KernelData *data)
+{
+ ccl_local unsigned int local_queue_atomics;
+ kernel_buffer_update((KernelGlobals*)kg, &local_queue_atomics);
+}
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
index ed64ae01aae..eb34f750881 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
@@ -18,9 +18,10 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_direct_lighting.h"
-#define KERNEL_NAME direct_lighting
-#define LOCALS_TYPE unsigned int
-#include "kernel/kernels/opencl/kernel_split_function.h"
-#undef KERNEL_NAME
-#undef LOCALS_TYPE
-
+__kernel void kernel_ocl_path_trace_direct_lighting(
+ ccl_global char *kg,
+ ccl_constant KernelData *data)
+{
+ ccl_local unsigned int local_queue_atomics;
+ kernel_direct_lighting((KernelGlobals*)kg, &local_queue_atomics);
+}
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
index 8afaa686e28..83ef5f5f3f2 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
@@ -18,7 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_do_volume.h"
-#define KERNEL_NAME do_volume
-#include "kernel/kernels/opencl/kernel_split_function.h"
-#undef KERNEL_NAME
-
+__kernel void kernel_ocl_path_trace_do_volume(
+ ccl_global char *kg,
+ ccl_constant KernelData *data)
+{
+ kernel_do_volume((KernelGlobals*)kg);
+}
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 9e1e57beba6..d071b39aa6f 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,9 +18,12 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h"
-#define KERNEL_NAME holdout_emission_blurring_pathtermination_ao
-#define LOCALS_TYPE BackgroundAOLocals
-#include "kernel/kernels/opencl/kernel_split_function.h"
-#undef KERNEL_NAME
-#undef LOCALS_TYPE
-
+__kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
+ ccl_global char *kg,
+ ccl_constant KernelData *data)
+{
+ ccl_local BackgroundAOLocals locals;
+ kernel_holdout_emission_blurring_pathtermination_ao(
+ (KernelGlobals*)kg,
+ &locals);
+}
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
index 192d01444ba..8c213ff5cb2 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
@@ -18,7 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_indirect_background.h"
-#define KERNEL_NAME indirect_background
-#include "kernel/kernels/opencl/kernel_split_function.h"
-#undef KERNEL_NAME
-
+__kernel void kernel_ocl_path_trace_indirect_background(
+ ccl_global char *kg,
+ ccl_constant KernelData *data)
+{
+ kernel_indirect_background((KernelGlobals*)kg);
+}
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
index 84938b889e5..998ebc4c0c3 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
@@ -18,7 +18,9 @@
#include "kernel/split/kernel_split_common.h"
#include "kernel/split/kernel_indirect_subsurface.h"
-#define KERNEL_NAME indirect_subsurface
-#include "kernel/kernels/opencl/kernel_split_function.h"
-#undef KERNEL_NAME
-
+__kernel void kernel_ocl_path_trace_indirect_subsurface(
+ ccl_global char *kg,
+ ccl_constant KernelData *data)
+{
+ kernel_indirect_subsurface((KernelGlobals*)kg);
+}
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
index c314dc96c33..822d2
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list