Cr8-xyz/blender
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Cycles: Optionally output luminance from the shader evaluation kernel

Browse Source 
This makes it possible to move some parts of evaluation from host to the device
and hopefully reduce memory usage by avoid having full RGBA buffer on the host.

Reviewers: juicyfruit, lukasstockner97, brecht

Reviewed By: lukasstockner97, brecht

Differential Revision: https://developer.blender.org/D1702
This commit is contained in:
Sergey Sharybin
2015-12-30 19:04:01 +05:00
parent c8a551bf13
commit 3918c8b9a5
13 changed files with 91 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
intern/cycles/device/device_cpu.cpp
View File

@@ -343,7 +343,7 @@ public:
#ifdef WITH_OSL
OSLShader::thread_init(&kg, &kernel_globals, &osl_globals);
#endif
void(*shader_kernel)(KernelGlobals*, uint4*, float4*, int, int, int, int);
void(*shader_kernel)(KernelGlobals*, uint4*, float4*, float*, int, int, int, int);
#ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX2
if(system_cpu_support_avx2())
@@ -374,8 +374,14 @@ public:
for(int sample = 0; sample < task.num_samples; sample++) {
for(int x = task.shader_x; x < task.shader_x + task.shader_w; x++)
shader_kernel(&kg, (uint4*)task.shader_input, (float4*)task.shader_output,
task.shader_eval_type, x, task.offset, sample);
shader_kernel(&kg,
(uint4*)task.shader_input,
(float4*)task.shader_output,
(float*)task.shader_output_luma,
task.shader_eval_type,
x,
task.offset,
sample);
if(task.get_cancel() || task_pool.canceled())
break;

20
intern/cycles/device/device_cuda.cpp
View File

@@ -726,6 +726,7 @@ public:
CUfunction cuShader;
CUdeviceptr d_input = cuda_device_ptr(task.shader_input);
CUdeviceptr d_output = cuda_device_ptr(task.shader_output);
CUdeviceptr d_output_luma = cuda_device_ptr(task.shader_output_luma);
/* get kernel function */
if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
@@ -747,13 +748,18 @@ public:
int shader_w = min(shader_chunk_size, end - shader_x);
/* pass in parameters */
void *args[] = {&d_input,
&d_output,
&task.shader_eval_type,
&shader_x,
&shader_w,
&offset,
&sample};
void *args[8];
int arg = 0;
args[arg++] = &d_input;
args[arg++] = &d_output;
if(task.shader_eval_type < SHADER_EVAL_BAKE) {
args[arg++] = &d_output_luma;
}
args[arg++] = &task.shader_eval_type;
args[arg++] = &shader_x;
args[arg++] = &shader_w;
args[arg++] = &offset;
args[arg++] = &sample;
/* launch kernel */
int threads_per_block;

1
intern/cycles/device/device_multi.cpp
View File

@@ -316,6 +316,7 @@ public:
if(task.rgba_half) subtask.rgba_half = sub.ptr_map[task.rgba_half];
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];
if(task.shader_output_luma) subtask.shader_output_luma = sub.ptr_map[task.shader_output_luma];
sub.device->task_add(subtask);
}

3
intern/cycles/device/device_network.cpp
View File

@@ -648,6 +648,9 @@ protected:
if(task.shader_output)
task.shader_output = device_ptr_from_client_pointer(task.shader_output);
if(task.shader_output)luma)
task.shader_output_luma = device_ptr_from_client_pointer(task.shader_output_luma);
task.acquire_tile = function_bind(&DeviceServer::task_acquire_tile, this, _1, _2);
task.release_tile = function_bind(&DeviceServer::task_release_tile, this, _1);

4
intern/cycles/device/device_network.h
View File

@@ -132,7 +132,7 @@ public:
archive & type & task.x & task.y & task.w & task.h;
archive & task.rgba_byte & task.rgba_half & task.buffer & task.sample & task.num_samples;
archive & task.offset & task.stride;
archive & task.shader_input & task.shader_output & task.shader_eval_type;
archive & task.shader_input & task.shader_output & task.shader_output_luma & task.shader_eval_type;
archive & task.shader_x & task.shader_w;
archive & task.need_finish_queue;
}
@@ -291,7 +291,7 @@ public:
*archive & type & task.x & task.y & task.w & task.h;
*archive & task.rgba_byte & task.rgba_half & task.buffer & task.sample & task.num_samples;
*archive & task.offset & task.stride;
*archive & task.shader_input & task.shader_output & task.shader_eval_type;
*archive & task.shader_input & task.shader_output & task.shader_output_luma & task.shader_eval_type;
*archive & task.shader_x & task.shader_w;
*archive & task.need_finish_queue;

7
intern/cycles/device/device_opencl.cpp
View File

@@ -1304,6 +1304,7 @@ public:
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_mem d_output_luma = CL_MEM_PTR(task.shader_output_luma);
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;
@@ -1330,6 +1331,12 @@ public:
d_input,
d_output);
if(task.shader_eval_type < SHADER_EVAL_BAKE) {
start_arg_index += kernel_set_args(kernel,
start_arg_index,
d_output_luma);
}
#define KERNEL_TEX(type, ttype, name) \
set_kernel_arg_mem(kernel, &start_arg_index, #name);
#include "kernel_textures.h"

2
intern/cycles/device/device_task.cpp
View File

@@ -29,7 +29,7 @@ CCL_NAMESPACE_BEGIN
DeviceTask::DeviceTask(Type type_)
: type(type_), x(0), y(0), w(0), h(0), rgba_byte(0), rgba_half(0), buffer(0),
sample(0), num_samples(1),
shader_input(0), shader_output(0),
shader_input(0), shader_output(0), shader_output_luma(0),
shader_eval_type(0), shader_x(0), shader_w(0)
{
last_update_time = time_dt();

2
intern/cycles/device/device_task.h
View File

@@ -46,7 +46,7 @@ public:
int offset, stride;
device_ptr shader_input;
device_ptr shader_output;
device_ptr shader_output, shader_output_luma;
int shader_eval_type;
int shader_x, shader_w;

24
intern/cycles/kernel/kernel_bake.h
View File

@@ -453,7 +453,13 @@ ccl_device void kernel_bake_evaluate(KernelGlobals *kg, ccl_global uint4 *input,
output[i] += make_float4(out.x, out.y, out.z, 1.0f) * output_fac;
}
ccl_device void kernel_shader_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_global float4 *output, ShaderEvalType type, int i, int sample)
ccl_device void kernel_shader_evaluate(KernelGlobals *kg,
ccl_global uint4 *input,
ccl_global float4 *output,
ccl_global float *output_luma,
ShaderEvalType type,
int i,
int sample)
{
ShaderData sd;
uint4 in = input[i];
@@ -500,10 +506,22 @@ ccl_device void kernel_shader_evaluate(KernelGlobals *kg, ccl_global uint4 *inpu
}
/* write output */
if(sample == 0)
if(sample == 0) {
if(output != NULL) {
output[i] = make_float4(out.x, out.y, out.z, 0.0f);
else
}
if(output_luma != NULL) {
output_luma[i] = average(out);
}
}
else {
if(output != NULL) {
output[i] += make_float4(out.x, out.y, out.z, 0.0f);
}
if(output_luma != NULL) {
output_luma[i] += average(out);
}
}
}
CCL_NAMESPACE_END

1
intern/cycles/kernel/kernels/cpu/kernel_cpu.h
View File

@@ -42,6 +42,7 @@ void KERNEL_FUNCTION_FULL_NAME(convert_to_half_float)(KernelGlobals *kg,
void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
uint4 *input,
float4 *output,
float *output_luma,
int type,
int i,
int offset,

3
intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
View File

@@ -99,12 +99,14 @@ void KERNEL_FUNCTION_FULL_NAME(convert_to_half_float)(KernelGlobals *kg,
void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
uint4 *input,
float4 *output,
float *output_luma,
int type,
int i,
int offset,
int sample)
{
if(type >= SHADER_EVAL_BAKE) {
kernel_assert(output_luma == NULL);
kernel_bake_evaluate(kg,
input,
output,
@@ -117,6 +119,7 @@ void KERNEL_FUNCTION_FULL_NAME(shader)(KernelGlobals *kg,
kernel_shader_evaluate(kg,
input,
output,
output_luma,
(ShaderEvalType)type,
i,
sample);

20
intern/cycles/kernel/kernels/cuda/kernel.cu
View File

@@ -159,12 +159,26 @@ kernel_cuda_convert_to_half_float(uchar4 *rgba, float *buffer, float sample_scal
extern "C" __global__ void
CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
kernel_cuda_shader(uint4 *input, float4 *output, int type, int sx, int sw, int offset, int sample)
kernel_cuda_shader(uint4 *input,
float4 *output,
float *output_luma,
int type,
int sx,
int sw,
int offset,
int sample)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
if(x < sx + sw)
kernel_shader_evaluate(NULL, input, output, (ShaderEvalType)type, x, sample);
if(x < sx + sw) {
kernel_shader_evaluate(NULL,
input,
output,
output_luma,
(ShaderEvalType)type,
x,
sample);
}
}
extern "C" __global__ void

12
intern/cycles/kernel/kernels/opencl/kernel.cl
View File

@@ -61,6 +61,7 @@ __kernel void kernel_ocl_shader(
ccl_constant KernelData *data,
ccl_global uint4 *input,
ccl_global float4 *output,
ccl_global float *output_luma,
#define KERNEL_TEX(type, ttype, name) \
ccl_global type *name,
@@ -78,8 +79,15 @@ __kernel void kernel_ocl_shader(
int x = sx + get_global_id(0);
if(x < sx + sw)
kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, x, sample);
if(x < sx + sw) {
kernel_shader_evaluate(kg,
input,
output,
output_luma,
(ShaderEvalType)type,
x,
sample);
}
}
__kernel void kernel_ocl_bake(