Revision: 43038
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=43038
Author: blendix
Date: 2011-12-31 15:18:13 +0000 (Sat, 31 Dec 2011)
Log Message:
-----------
Cycles code refactoring: change displace kernel into more generic shader
evaluate kernel, added background shader evaluate.
Modified Paths:
--------------
trunk/blender/intern/cycles/device/device.cpp
trunk/blender/intern/cycles/device/device.h
trunk/blender/intern/cycles/device/device_cpu.cpp
trunk/blender/intern/cycles/device/device_cuda.cpp
trunk/blender/intern/cycles/device/device_multi.cpp
trunk/blender/intern/cycles/kernel/kernel.cl
trunk/blender/intern/cycles/kernel/kernel.cpp
trunk/blender/intern/cycles/kernel/kernel.cu
trunk/blender/intern/cycles/kernel/kernel.h
trunk/blender/intern/cycles/kernel/kernel_displace.h
trunk/blender/intern/cycles/kernel/kernel_optimized.cpp
trunk/blender/intern/cycles/kernel/kernel_types.h
trunk/blender/intern/cycles/render/mesh_displace.cpp
Modified: trunk/blender/intern/cycles/device/device.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device.cpp 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/device/device.cpp 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -38,7 +38,8 @@
DeviceTask::DeviceTask(Type type_)
: type(type_), x(0), y(0), w(0), h(0), rng_state(0), rgba(0), buffer(0),
sample(0), resolution(0),
- displace_input(0), displace_offset(0), displace_x(0), displace_w(0)
+ shader_input(0), shader_output(0),
+ shader_eval_type(0), shader_x(0), shader_w(0)
{
}
@@ -46,8 +47,8 @@
{
int num;
- if(type == DISPLACE) {
- num = (displace_w + max_size - 1)/max_size;
+ if(type == SHADER) {
+ num = (shader_w + max_size - 1)/max_size;
}
else {
max_size = max(1, max_size/w);
@@ -68,17 +69,17 @@
void DeviceTask::split(list<DeviceTask>& tasks, int num)
{
- if(type == DISPLACE) {
- num = min(displace_w, num);
+ if(type == SHADER) {
+ num = min(shader_w, num);
for(int i = 0; i < num; i++) {
- int tx = displace_x + (displace_w/num)*i;
- int tw = (i == num-1)? displace_w - i*(displace_w/num):
displace_w/num;
+ int tx = shader_x + (shader_w/num)*i;
+ int tw = (i == num-1)? shader_w - i*(shader_w/num):
shader_w/num;
DeviceTask task = *this;
- task.displace_x = tx;
- task.displace_w = tw;
+ task.shader_x = tx;
+ task.shader_w = tw;
tasks.push_back(task);
}
Modified: trunk/blender/intern/cycles/device/device.h
===================================================================
--- trunk/blender/intern/cycles/device/device.h 2011-12-31 15:10:38 UTC (rev
43037)
+++ trunk/blender/intern/cycles/device/device.h 2011-12-31 15:18:13 UTC (rev
43038)
@@ -52,7 +52,7 @@
class DeviceTask {
public:
- typedef enum { PATH_TRACE, TONEMAP, DISPLACE } Type;
+ typedef enum { PATH_TRACE, TONEMAP, SHADER } Type;
Type type;
int x, y, w, h;
@@ -63,9 +63,10 @@
int resolution;
int offset, stride;
- device_ptr displace_input;
- device_ptr displace_offset;
- int displace_x, displace_w;
+ device_ptr shader_input;
+ device_ptr shader_output;
+ int shader_eval_type;
+ int shader_x, shader_w;
DeviceTask(Type type = PATH_TRACE);
Modified: trunk/blender/intern/cycles/device/device_cpu.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_cpu.cpp 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/device/device_cpu.cpp 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -141,8 +141,8 @@
thread_path_trace(task);
else if(task.type == DeviceTask::TONEMAP)
thread_tonemap(task);
- else if(task.type == DeviceTask::DISPLACE)
- thread_displace(task);
+ else if(task.type == DeviceTask::SHADER)
+ thread_shader(task);
tasks.worker_done();
}
@@ -207,7 +207,7 @@
}
}
- void thread_displace(DeviceTask& task)
+ void thread_shader(DeviceTask& task)
{
#ifdef WITH_OSL
if(kernel_osl_use(kg))
@@ -216,8 +216,8 @@
#ifdef WITH_OPTIMIZED_KERNEL
if(system_cpu_support_optimized()) {
- for(int x = task.displace_x; x < task.displace_x +
task.displace_w; x++) {
- kernel_cpu_optimized_displace(kg,
(uint4*)task.displace_input, (float3*)task.displace_offset, x);
+ for(int x = task.shader_x; x < task.shader_x +
task.shader_w; x++) {
+ kernel_cpu_optimized_shader(kg,
(uint4*)task.shader_input, (float3*)task.shader_output, task.shader_eval_type,
x);
if(tasks.worker_cancel())
break;
@@ -226,8 +226,8 @@
else
#endif
{
- for(int x = task.displace_x; x < task.displace_x +
task.displace_w; x++) {
- kernel_cpu_displace(kg,
(uint4*)task.displace_input, (float3*)task.displace_offset, x);
+ for(int x = task.shader_x; x < task.shader_x +
task.shader_w; x++) {
+ kernel_cpu_shader(kg,
(uint4*)task.shader_input, (float3*)task.shader_output, task.shader_eval_type,
x);
if(tasks.worker_cancel())
break;
Modified: trunk/blender/intern/cycles/device/device_cuda.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_cuda.cpp 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/device/device_cuda.cpp 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -615,16 +615,16 @@
cuda_pop_context();
}
- void displace(DeviceTask& task)
+ void shader(DeviceTask& task)
{
cuda_push_context();
CUfunction cuDisplace;
- CUdeviceptr d_input = cuda_device_ptr(task.displace_input);
- CUdeviceptr d_offset = cuda_device_ptr(task.displace_offset);
+ CUdeviceptr d_input = cuda_device_ptr(task.shader_input);
+ CUdeviceptr d_offset = cuda_device_ptr(task.shader_output);
/* get kernel function */
- cuda_assert(cuModuleGetFunction(&cuDisplace, cuModule,
"kernel_cuda_displace"))
+ cuda_assert(cuModuleGetFunction(&cuDisplace, cuModule,
"kernel_cuda_shader"))
/* pass in parameters */
int offset = 0;
@@ -635,12 +635,15 @@
cuda_assert(cuParamSetv(cuDisplace, offset, &d_offset,
sizeof(d_offset)))
offset += sizeof(d_offset);
- int displace_x = task.displace_x;
- offset = cuda_align_up(offset, __alignof(displace_x));
+ int shader_eval_type = task.shader_eval_type;
+ offset = cuda_align_up(offset, __alignof(shader_eval_type));
- cuda_assert(cuParamSeti(cuDisplace, offset, task.displace_x))
- offset += sizeof(task.displace_x);
+ cuda_assert(cuParamSeti(cuDisplace, offset,
task.shader_eval_type))
+ offset += sizeof(task.shader_eval_type);
+ cuda_assert(cuParamSeti(cuDisplace, offset, task.shader_x))
+ offset += sizeof(task.shader_x);
+
cuda_assert(cuParamSetSize(cuDisplace, offset))
/* launch kernel: todo find optimal size, cache config for
fermi */
@@ -649,7 +652,7 @@
#else
int xthreads = 8;
#endif
- int xblocks = (task.displace_w + xthreads - 1)/xthreads;
+ int xblocks = (task.shader_w + xthreads - 1)/xthreads;
cuda_assert(cuFuncSetCacheConfig(cuDisplace,
CU_FUNC_CACHE_PREFER_L1))
cuda_assert(cuFuncSetBlockShape(cuDisplace, xthreads, 1, 1))
@@ -828,8 +831,8 @@
tonemap(task);
else if(task.type == DeviceTask::PATH_TRACE)
path_trace(task);
- else if(task.type == DeviceTask::DISPLACE)
- displace(task);
+ else if(task.type == DeviceTask::SHADER)
+ shader(task);
}
void task_wait()
Modified: trunk/blender/intern/cycles/device/device_multi.cpp
===================================================================
--- trunk/blender/intern/cycles/device/device_multi.cpp 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/device/device_multi.cpp 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -306,8 +306,8 @@
if(task.buffer) subtask.buffer =
sub.ptr_map[task.buffer];
if(task.rng_state) subtask.rng_state =
sub.ptr_map[task.rng_state];
if(task.rgba) subtask.rgba =
sub.ptr_map[task.rgba];
- if(task.displace_input) subtask.displace_input
= sub.ptr_map[task.displace_input];
- if(task.displace_offset)
subtask.displace_offset = sub.ptr_map[task.displace_offset];
+ if(task.shader_input) subtask.shader_input =
sub.ptr_map[task.shader_input];
+ if(task.shader_output) subtask.shader_output =
sub.ptr_map[task.shader_output];
sub.device->task_add(subtask);
}
Modified: trunk/blender/intern/cycles/kernel/kernel.cl
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cl 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/kernel/kernel.cl 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -80,10 +80,10 @@
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y,
offset, stride);
}
-/*__kernel void kernel_ocl_displace(__global uint4 *input, __global float3
*offset, int sx)
+/*__kernel void kernel_ocl_shader(__global uint4 *input, __global float3
*output, int type, int sx)
{
int x = sx + get_global_id(0);
- kernel_displace(input, offset, x);
+ kernel_shader_evaluate(input, output, (ShaderEvalType)type, x);
}*/
Modified: trunk/blender/intern/cycles/kernel/kernel.cpp
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cpp 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/kernel/kernel.cpp 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -216,11 +216,11 @@
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset,
stride);
}
-/* Displacement */
+/* Shader Evaluation */
-void kernel_cpu_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int
i)
+void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float3 *output, int
type, int i)
{
- kernel_displace(kg, input, offset, i);
+ kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i);
}
CCL_NAMESPACE_END
Modified: trunk/blender/intern/cycles/kernel/kernel.cu
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.cu 2011-12-31 15:10:38 UTC
(rev 43037)
+++ trunk/blender/intern/cycles/kernel/kernel.cu 2011-12-31 15:18:13 UTC
(rev 43038)
@@ -44,10 +44,10 @@
kernel_film_tonemap(NULL, rgba, buffer, sample, resolution, x,
y, offset, stride);
}
-extern "C" __global__ void kernel_cuda_displace(uint4 *input, float3 *offset,
int sx)
+extern "C" __global__ void kernel_cuda_shader(uint4 *input, float3 *output,
int type, int sx)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
- kernel_displace(NULL, input, offset, x);
+ kernel_shader_evaluate(NULL, input, output, (ShaderEvalType)type, x);
}
Modified: trunk/blender/intern/cycles/kernel/kernel.h
===================================================================
--- trunk/blender/intern/cycles/kernel/kernel.h 2011-12-31 15:10:38 UTC (rev
43037)
+++ trunk/blender/intern/cycles/kernel/kernel.h 2011-12-31 15:18:13 UTC (rev
43038)
@@ -40,14 +40,16 @@
int sample, int x, int y, int offset, int stride);
void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
int sample, int resolution, int x, int y, int offset, int stride);
-void kernel_cpu_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int
i);
+void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float3 *output,
+ int type, int i);
#ifdef WITH_OPTIMIZED_KERNEL
void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer,
unsigned int *rng_state,
int sample, int x, int y, int offset, int stride);
void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4
*buffer,
@@ Diff output truncated at 10240 characters. @@
_______________________________________________
Bf-blender-cvs mailing list
[email protected]
http://lists.blender.org/mailman/listinfo/bf-blender-cvs
