[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [57636] trunk/blender/intern/cycles: Cycles OpenCL: make displacement and world importance sampling work.
Brecht Van Lommel
brechtvanlommel at pandora.be
Fri Jun 21 15:05:08 CEST 2013
Revision: 57636
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=57636
Author: blendix
Date: 2013-06-21 13:05:08 +0000 (Fri, 21 Jun 2013)
Log Message:
-----------
Cycles OpenCL: make displacement and world importance sampling work.
Modified Paths:
--------------
trunk/blender/intern/cycles/device/device_cuda.cpp
trunk/blender/intern/cycles/device/device_opencl.cpp
trunk/blender/intern/cycles/kernel/kernel.cl
trunk/blender/intern/cycles/kernel/kernel_displace.h
trunk/blender/intern/cycles/kernel/kernel_types.h
trunk/blender/intern/cycles/render/light.cpp
Modified: trunk/blender/intern/cycles/device/device_cuda.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_cuda.cpp 2013-06-21 12:57:25 UTC (rev 57635)
+++ trunk/blender/intern/cycles/device/device_cuda.cpp 2013-06-21 13:05:08 UTC (rev 57636)
@@ -704,7 +704,7 @@
CUfunction cuDisplace;
CUdeviceptr d_input = cuda_device_ptr(task.shader_input);
- CUdeviceptr d_offset = cuda_device_ptr(task.shader_output);
+ CUdeviceptr d_output = cuda_device_ptr(task.shader_output);
/* get kernel function */
cuda_assert(cuModuleGetFunction(&cuDisplace, cuModule, "kernel_cuda_shader"))
@@ -715,8 +715,8 @@
cuda_assert(cuParamSetv(cuDisplace, offset, &d_input, sizeof(d_input)))
offset += sizeof(d_input);
- cuda_assert(cuParamSetv(cuDisplace, offset, &d_offset, sizeof(d_offset)))
- offset += sizeof(d_offset);
+ cuda_assert(cuParamSetv(cuDisplace, offset, &d_output, sizeof(d_output)))
+ offset += sizeof(d_output);
int shader_eval_type = task.shader_eval_type;
offset = align_up(offset, __alignof(shader_eval_type));
Modified: trunk/blender/intern/cycles/device/device_opencl.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_opencl.cpp 2013-06-21 12:57:25 UTC (rev 57635)
+++ trunk/blender/intern/cycles/device/device_opencl.cpp 2013-06-21 13:05:08 UTC (rev 57636)
@@ -318,6 +318,7 @@
cl_program cpProgram;
cl_kernel ckPathTraceKernel;
cl_kernel ckFilmConvertKernel;
+ cl_kernel ckShaderKernel;
cl_int ciErr;
typedef map<string, device_vector<uchar>*> ConstMemMap;
@@ -427,6 +428,7 @@
cpProgram = NULL;
ckPathTraceKernel = NULL;
ckFilmConvertKernel = NULL;
+ ckShaderKernel = NULL;
null_mem = 0;
device_initialized = false;
@@ -760,6 +762,10 @@
if(opencl_error(ciErr))
return false;
+ ckShaderKernel = clCreateKernel(cpProgram, "kernel_ocl_shader", &ciErr);
+ if(opencl_error(ciErr))
+ return false;
+
return true;
}
@@ -1009,11 +1015,45 @@
enqueue_kernel(ckFilmConvertKernel, d_w, d_h);
}
+ void shader(DeviceTask& task)
+ {
+ /* cast arguments to cl types */
+ cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
+ cl_mem d_input = CL_MEM_PTR(task.shader_input);
+ cl_mem d_output = CL_MEM_PTR(task.shader_output);
+ cl_int d_shader_eval_type = task.shader_eval_type;
+ cl_int d_shader_x = task.shader_x;
+ cl_int d_shader_w = task.shader_w;
+
+ /* sample arguments */
+ cl_uint narg = 0;
+ ciErr = 0;
+
+ ciErr |= clSetKernelArg(ckShaderKernel, narg++, sizeof(d_data), (void*)&d_data);
+ ciErr |= clSetKernelArg(ckShaderKernel, narg++, sizeof(d_input), (void*)&d_input);
+ ciErr |= clSetKernelArg(ckShaderKernel, narg++, sizeof(d_output), (void*)&d_output);
+
+#define KERNEL_TEX(type, ttype, name) \
+ ciErr |= set_kernel_arg_mem(ckShaderKernel, &narg, #name);
+#include "kernel_textures.h"
+
+ ciErr |= clSetKernelArg(ckShaderKernel, narg++, sizeof(d_shader_eval_type), (void*)&d_shader_eval_type);
+ ciErr |= clSetKernelArg(ckShaderKernel, narg++, sizeof(d_shader_x), (void*)&d_shader_x);
+ ciErr |= clSetKernelArg(ckShaderKernel, narg++, sizeof(d_shader_w), (void*)&d_shader_w);
+
+ opencl_assert(ciErr);
+
+ enqueue_kernel(ckShaderKernel, task.shader_w, 1);
+ }
+
void thread_run(DeviceTask *task)
{
if(task->type == DeviceTask::TONEMAP) {
tonemap(*task, task->buffer, task->rgba);
}
+ else if(task->type == DeviceTask::SHADER) {
+ shader(*task);
+ }
else if(task->type == DeviceTask::PATH_TRACE) {
RenderTile tile;
Modified: trunk/blender/intern/cycles/kernel/kernel.cl
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cl 2013-06-21 12:57:25 UTC (rev 57635)
+++ trunk/blender/intern/cycles/kernel/kernel.cl 2013-06-21 13:05:08 UTC (rev 57636)
@@ -25,6 +25,7 @@
#include "kernel_film.h"
#include "kernel_path.h"
+#include "kernel_displace.h"
__kernel void kernel_ocl_path_trace(
__constant KernelData *data,
@@ -80,10 +81,28 @@
kernel_film_tonemap(kg, rgba, buffer, sample, x, y, offset, stride);
}
-/*__kernel void kernel_ocl_shader(__global uint4 *input, __global float *output, int type, int sx)
+__kernel void kernel_ocl_shader(
+ __constant KernelData *data,
+ __global uint4 *input,
+ __global float4 *output,
+
+#define KERNEL_TEX(type, ttype, name) \
+ __global type *name,
+#include "kernel_textures.h"
+
+ int type, int sx, int sw)
{
+ KernelGlobals kglobals, *kg = &kglobals;
+
+ kg->data = data;
+
+#define KERNEL_TEX(type, ttype, name) \
+ kg->name = name;
+#include "kernel_textures.h"
+
int x = sx + get_global_id(0);
- kernel_shader_evaluate(input, output, (ShaderEvalType)type, x);
-}*/
+ if(x < sx + sw)
+ kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, x);
+}
Modified: trunk/blender/intern/cycles/kernel/kernel_displace.h
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel_displace.h 2013-06-21 12:57:25 UTC (rev 57635)
+++ trunk/blender/intern/cycles/kernel/kernel_displace.h 2013-06-21 13:05:08 UTC (rev 57636)
@@ -18,7 +18,7 @@
CCL_NAMESPACE_BEGIN
-__device void kernel_shader_evaluate(KernelGlobals *kg, uint4 *input, float4 *output, ShaderEvalType type, int i)
+__device void kernel_shader_evaluate(KernelGlobals *kg, __global uint4 *input, __global float4 *output, ShaderEvalType type, int i)
{
ShaderData sd;
uint4 in = input[i];
Modified: trunk/blender/intern/cycles/kernel/kernel_types.h
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel_types.h 2013-06-21 12:57:25 UTC (rev 57635)
+++ trunk/blender/intern/cycles/kernel/kernel_types.h 2013-06-21 13:05:08 UTC (rev 57636)
@@ -150,10 +150,10 @@
/* Shader Evaluation */
-enum ShaderEvalType {
+typedef enum ShaderEvalType {
SHADER_EVAL_DISPLACE,
SHADER_EVAL_BACKGROUND
-};
+} ShaderEvalType;
/* Path Tracing
* note we need to keep the u/v pairs at even values */
Modified: trunk/blender/intern/cycles/render/light.cpp
===================================================================
--- trunk/blender/intern/cycles/render/light.cpp 2013-06-21 12:57:25 UTC (rev 57635)
+++ trunk/blender/intern/cycles/render/light.cpp 2013-06-21 13:05:08 UTC (rev 57636)
@@ -30,7 +30,7 @@
CCL_NAMESPACE_BEGIN
-static void dump_background_pixels(Device *device, DeviceScene *dscene, int res, vector<float3>& pixels)
+static void shade_background_pixels(Device *device, DeviceScene *dscene, int res, vector<float3>& pixels)
{
/* create input */
int width = res;
@@ -433,7 +433,7 @@
assert(res > 0);
vector<float3> pixels;
- dump_background_pixels(device, dscene, res, pixels);
+ shade_background_pixels(device, dscene, res, pixels);
if(progress.get_cancel())
return;
More information about the Bf-blender-cvs
mailing list