griefith/test2
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

svn merge ^/trunk/blender -r42761:42776

Browse Source
This commit is contained in:
Campbell Barton
2011-12-20 22:08:24 +00:00
parent deb3b7e282 2d1de2e78d
commit 1fef05084e
41 changed files with 389 additions and 207 deletions
Download Patch File Download Diff File Expand all files Collapse all files

19
intern/cycles/app/cycles_test.cpp
View File

@@ -82,10 +82,21 @@ static void session_print_status()
session_print(status);
}
static BufferParams session_buffer_params()
{
BufferParams buffer_params;
buffer_params.width = options.width;
buffer_params.height = options.height;
buffer_params.full_width = options.width;
buffer_params.full_height = options.height;
return buffer_params;
}
static void session_init()
{
options.session = new Session(options.session_params);
options.session->reset(options.width, options.height, options.session_params.samples);
options.session->reset(session_buffer_params(), options.session_params.samples);
options.session->scene = options.scene;
if(options.session_params.background && !options.quiet)
@@ -151,7 +162,7 @@ static void display_info(Progress& progress)
static void display()
{
options.session->draw(options.width, options.height);
options.session->draw(session_buffer_params());
display_info(options.session->progress);
}
@@ -162,13 +173,13 @@ static void resize(int width, int height)
options.height= height;
if(options.session)
options.session->reset(options.width, options.height, options.session_params.samples);
options.session->reset(session_buffer_params(), options.session_params.samples);
}
void keyboard(unsigned char key)
{
if(key == 'r')
options.session->reset(options.width, options.height, options.session_params.samples);
options.session->reset(session_buffer_params(), options.session_params.samples);
else if(key == 27) // escape
options.session->progress.set_cancel("Cancelled");
}

6
intern/cycles/blender/addon/engine.py
View File

@@ -62,11 +62,7 @@ def render(engine):
def update(engine, data, scene):
import bcycles
if scene.render.use_border:
engine.report({'ERROR'}, "Border rendering not supported yet")
free(engine)
else:
bcycles.sync(engine.session)
bcycles.sync(engine.session)
def draw(engine, region, v3d, rv3d):

24
intern/cycles/blender/blender_camera.cpp
View File

@@ -287,5 +287,29 @@ void BlenderSync::sync_view(BL::SpaceView3D b_v3d, BL::RegionView3D b_rv3d, int
blender_camera_sync(scene->camera, &bcam, width, height);
}
BufferParams BlenderSync::get_buffer_params(BL::Scene b_scene, BL::RegionView3D b_rv3d, int width, int height)
{
BufferParams params;
params.full_width = width;
params.full_height = height;
/* border render */
BL::RenderSettings r = b_scene.render();
if(!b_rv3d && r.use_border()) {
params.full_x = r.border_min_x()*width;
params.full_y = r.border_min_y()*height;
params.width = (int)(r.border_max_x()*width) - params.full_x;
params.height = (int)(r.border_max_y()*height) - params.full_y;
}
else {
params.width = width;
params.height = height;
}
return params;
}
CCL_NAMESPACE_END

23
intern/cycles/blender/blender_session.cpp
View File

@@ -100,7 +100,9 @@ void BlenderSession::create_session()
session->set_pause(BlenderSync::get_session_pause(b_scene, background));
/* start rendering */
session->reset(width, height, session_params.samples);
BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
session->reset(buffer_params, session_params.samples);
session->start();
}
@@ -135,7 +137,10 @@ void BlenderSession::write_render_result()
if(!pixels)
return;
struct RenderResult *rrp = RE_engine_begin_result((RenderEngine*)b_engine.ptr.data, 0, 0, width, height);
BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
int w = buffer_params.width, h = buffer_params.height;
struct RenderResult *rrp = RE_engine_begin_result((RenderEngine*)b_engine.ptr.data, 0, 0, w, h);
PointerRNA rrptr;
RNA_pointer_create(NULL, &RNA_RenderResult, rrp, &rrptr);
BL::RenderResult rr(rrptr);
@@ -188,8 +193,10 @@ void BlenderSession::synchronize()
session->scene->mutex.unlock();
/* reset if needed */
if(scene->need_reset())
session->reset(width, height, session_params.samples);
if(scene->need_reset()) {
BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
session->reset(buffer_params, session_params.samples);
}
}
bool BlenderSession::draw(int w, int h)
@@ -225,7 +232,9 @@ bool BlenderSession::draw(int w, int h)
/* reset if requested */
if(reset) {
SessionParams session_params = BlenderSync::get_session_params(b_scene, background);
session->reset(width, height, session_params.samples);
BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
session->reset(buffer_params, session_params.samples);
}
}
@@ -233,7 +242,9 @@ bool BlenderSession::draw(int w, int h)
update_status_progress();
/* draw */
return !session->draw(width, height);
BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
return !session->draw(buffer_params);
}
void BlenderSession::get_status(string& status, string& substatus)

1
intern/cycles/blender/blender_sync.h
View File

@@ -62,6 +62,7 @@ public:
static SceneParams get_scene_params(BL::Scene b_scene, bool background);
static SessionParams get_session_params(BL::Scene b_scene, bool background);
static bool get_session_pause(BL::Scene b_scene, bool background);
static BufferParams get_buffer_params(BL::Scene b_scene, BL::RegionView3D b_rv3d, int width, int height);
private:
/* sync */

31
intern/cycles/device/device.cpp
View File

@@ -24,6 +24,7 @@
#include "util_cuda.h"
#include "util_debug.h"
#include "util_foreach.h"
#include "util_math.h"
#include "util_opencl.h"
#include "util_opengl.h"
@@ -41,7 +42,31 @@ DeviceTask::DeviceTask(Type type_)
{
}
void DeviceTask::split(ThreadQueue<DeviceTask>& tasks, int num)
void DeviceTask::split_max_size(list<DeviceTask>& tasks, int max_size)
{
int num;
if(type == DISPLACE) {
num = (displace_w + max_size - 1)/max_size;
}
else {
max_size = max(1, max_size/w);
num = (h + max_size - 1)/max_size;
}
split(tasks, num);
}
void DeviceTask::split(ThreadQueue<DeviceTask>& queue, int num)
{
list<DeviceTask> tasks;
split(tasks, num);
foreach(DeviceTask& task, tasks)
queue.push(task);
}
void DeviceTask::split(list<DeviceTask>& tasks, int num)
{
if(type == DISPLACE) {
num = min(displace_w, num);
@@ -55,7 +80,7 @@ void DeviceTask::split(ThreadQueue<DeviceTask>& tasks, int num)
task.displace_x = tx;
task.displace_w = tw;
tasks.push(task);
tasks.push_back(task);
}
}
else {
@@ -70,7 +95,7 @@ void DeviceTask::split(ThreadQueue<DeviceTask>& tasks, int num)
task.y = ty;
task.h = th;
tasks.push(task);
tasks.push_back(task);
}
}
}

5
intern/cycles/device/device.h
View File

@@ -23,6 +23,7 @@
#include "device_memory.h"
#include "util_list.h"
#include "util_string.h"
#include "util_thread.h"
#include "util_types.h"
@@ -60,13 +61,17 @@ public:
device_ptr buffer;
int sample;
int resolution;
int offset, stride;
device_ptr displace_input;
device_ptr displace_offset;
int displace_x, displace_w;
DeviceTask(Type type = PATH_TRACE);
void split(list<DeviceTask>& tasks, int num);
void split(ThreadQueue<DeviceTask>& tasks, int num);
void split_max_size(list<DeviceTask>& tasks, int max_size);
};
/* Device */

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

@@ -162,7 +162,8 @@ public:
if(system_cpu_support_optimized()) {
for(int y = task.y; y < task.y + task.h; y++) {
for(int x = task.x; x < task.x + task.w; x++)
kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state, task.sample, x, y);
kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
task.sample, x, y, task.offset, task.stride);
if(tasks.worker_cancel())
break;
@@ -173,7 +174,8 @@ public:
{
for(int y = task.y; y < task.y + task.h; y++) {
for(int x = task.x; x < task.x + task.w; x++)
kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state, task.sample, x, y);
kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
task.sample, x, y, task.offset, task.stride);
if(tasks.worker_cancel())
break;
@@ -192,14 +194,16 @@ public:
if(system_cpu_support_optimized()) {
for(int y = task.y; y < task.y + task.h; y++)
for(int x = task.x; x < task.x + task.w; x++)
kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer, task.sample, task.resolution, x, y);
kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
task.sample, task.resolution, x, y, task.offset, task.stride);
}
else
#endif
{
for(int y = task.y; y < task.y + task.h; y++)
for(int x = task.x; x < task.x + task.w; x++)
kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer, task.sample, task.resolution, x, y);
kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
task.sample, task.resolution, x, y, task.offset, task.stride);
}
}

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

@@ -221,7 +221,7 @@ public:
cuDeviceComputeCapability(&major, &minor, cuDevId);
if(major <= 1 && minor <= 2) {
cuda_error(string_printf("CUDA device supported only with shader model 1.3 or up, found %d.%d.", major, minor));
cuda_error(string_printf("CUDA device supported only with compute capability 1.3 or up, found %d.%d.", major, minor));
return false;
}
}
@@ -253,9 +253,9 @@ public:
#if defined(WITH_CUDA_BINARIES) && defined(_WIN32)
if(major <= 1 && minor <= 2)
cuda_error(string_printf("CUDA device supported only with shader model 1.3 or up, found %d.%d.", major, minor));
cuda_error(string_printf("CUDA device supported only compute capability 1.3 or up, found %d.%d.", major, minor));
else
cuda_error(string_printf("CUDA binary kernel for this graphics card shader model (%d.%d) not found.", major, minor));
cuda_error(string_printf("CUDA binary kernel for this graphics card compute capability (%d.%d) not found.", major, minor));
return "";
#else
/* if not, find CUDA compiler */
@@ -520,6 +520,12 @@ public:
cuda_assert(cuParamSeti(cuPathTrace, offset, task.h))
offset += sizeof(task.h);
cuda_assert(cuParamSeti(cuPathTrace, offset, task.offset))
offset += sizeof(task.offset);
cuda_assert(cuParamSeti(cuPathTrace, offset, task.stride))
offset += sizeof(task.stride);
cuda_assert(cuParamSetSize(cuPathTrace, offset))
/* launch kernel: todo find optimal size, cache config for fermi */
@@ -581,6 +587,12 @@ public:
cuda_assert(cuParamSeti(cuFilmConvert, offset, task.h))
offset += sizeof(task.h);
cuda_assert(cuParamSeti(cuFilmConvert, offset, task.offset))
offset += sizeof(task.offset);
cuda_assert(cuParamSeti(cuFilmConvert, offset, task.stride))
offset += sizeof(task.stride);
cuda_assert(cuParamSetSize(cuFilmConvert, offset))
/* launch kernel: todo find optimal size, cache config for fermi */

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

@@ -25,6 +25,7 @@
#include "device.h"
#include "device_intern.h"
#include "util_foreach.h"
#include "util_map.h"
#include "util_math.h"
#include "util_md5.h"
@@ -52,6 +53,7 @@ public:
map<string, device_memory*> mem_map;
device_ptr null_mem;
bool device_initialized;
string platform_name;
const char *opencl_error_string(cl_int err)
{
@@ -175,6 +177,10 @@ public:
if(opencl_error(ciErr))
return;
char name[256];
clGetPlatformInfo(cpPlatform, CL_PLATFORM_NAME, sizeof(name), &name, NULL);
platform_name = name;
cxContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr);
if(opencl_error(ciErr))
return;
@@ -277,14 +283,11 @@ public:
{
string build_options = " -cl-fast-relaxed-math ";
/* Full Shading only on NVIDIA cards at the moment */
char vendor[256];
clGetPlatformInfo(cpPlatform, CL_PLATFORM_NAME, sizeof(vendor), &vendor, NULL);
string name = vendor;
if(name == "NVIDIA CUDA")
build_options += "-D__KERNEL_SHADING__ -D__MULTI_CLOSURE__ ";
/* full shading only on NVIDIA cards at the moment */
if(platform_name == "NVIDIA CUDA")
build_options += "-D__KERNEL_SHADING__ -D__MULTI_CLOSURE__ -cl-nv-maxrregcount=24 -cl-nv-verbose ";
if(platform_name == "Apple")
build_options += " -D__CL_NO_FLOAT3__ ";
return build_options;
}
@@ -541,6 +544,8 @@ public:
cl_int d_w = task.w;
cl_int d_h = task.h;
cl_int d_sample = task.sample;
cl_int d_offset = task.offset;
cl_int d_stride = task.stride;
/* sample arguments */
int narg = 0;
@@ -559,6 +564,8 @@ public:
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_y), (void*)&d_y);
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_w), (void*)&d_w);
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_h), (void*)&d_h);
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_offset), (void*)&d_offset);
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_stride), (void*)&d_stride);
opencl_assert(ciErr);
@@ -611,6 +618,8 @@ public:
cl_int d_h = task.h;
cl_int d_sample = task.sample;
cl_int d_resolution = task.resolution;
cl_int d_offset = task.offset;
cl_int d_stride = task.stride;
/* sample arguments */
int narg = 0;
@@ -630,6 +639,8 @@ public:
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_y), (void*)&d_y);
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_w), (void*)&d_w);
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_h), (void*)&d_h);
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_offset), (void*)&d_offset);
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_stride), (void*)&d_stride);
opencl_assert(ciErr);
@@ -649,12 +660,24 @@ public:
opencl_assert(clFinish(cqCommandQueue));
}
void task_add(DeviceTask& task)
void task_add(DeviceTask& maintask)
{
if(task.type == DeviceTask::TONEMAP)
tonemap(task);
else if(task.type == DeviceTask::PATH_TRACE)
path_trace(task);
list<DeviceTask> tasks;
/* arbitrary limit to work around apple ATI opencl issue */
if(platform_name == "Apple")
maintask.split_max_size(tasks, 76800);
else
tasks.push_back(maintask);
DeviceTask task;
foreach(DeviceTask& task, tasks) {
if(task.type == DeviceTask::TONEMAP)
tonemap(task);
else if(task.type == DeviceTask::PATH_TRACE)
path_trace(task);
}
}
void task_wait()

2
intern/cycles/kernel/CMakeLists.txt
View File

@@ -143,7 +143,7 @@ endif()
#set(KERNEL_PREPROCESSED ${CMAKE_CURRENT_BINARY_DIR}/kernel_preprocessed.cl)
#add_custom_command(
# OUTPUT ${KERNEL_PREPROCESSED}
# COMMAND gcc -x c++ -E ${CMAKE_CURRENT_SOURCE_DIR}/kernel.cl -I ${CMAKE_CURRENT_SOURCE_DIR}/../util/ -DCCL_NAMESPACE_BEGIN= -DCCL_NAMESPACE_END= -DWITH_OPENCL -o ${KERNEL_PREPROCESSED}
# COMMAND gcc -x c++ -E ${CMAKE_CURRENT_SOURCE_DIR}/kernel.cl -I ${CMAKE_CURRENT_SOURCE_DIR}/../util/ -DCCL_NAMESPACE_BEGIN= -DCCL_NAMESPACE_END= -o ${KERNEL_PREPROCESSED}
# DEPENDS ${SRC_KERNEL} ${SRC_UTIL_HEADERS})
#add_custom_target(cycles_kernel_preprocess ALL DEPENDS ${KERNEL_PREPROCESSED})
#delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${KERNEL_PREPROCESSED}" ${CYCLES_INSTALL_PATH}/kernel)

8
intern/cycles/kernel/kernel.cl
View File

@@ -36,7 +36,7 @@ __kernel void kernel_ocl_path_trace(
#include "kernel_textures.h"
int sample,
int sx, int sy, int sw, int sh)
int sx, int sy, int sw, int sh, int offset, int stride)
{
KernelGlobals kglobals, *kg = &kglobals;
@@ -50,7 +50,7 @@ __kernel void kernel_ocl_path_trace(
int y = sy + get_global_id(1);
if(x < sx + sw && y < sy + sh)
kernel_path_trace(kg, buffer, rng_state, sample, x, y);
kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
}
__kernel void kernel_ocl_tonemap(
@@ -63,7 +63,7 @@ __kernel void kernel_ocl_tonemap(
#include "kernel_textures.h"
int sample, int resolution,
int sx, int sy, int sw, int sh)
int sx, int sy, int sw, int sh, int offset, int stride)
{
KernelGlobals kglobals, *kg = &kglobals;
@@ -77,7 +77,7 @@ __kernel void kernel_ocl_tonemap(
int y = sy + get_global_id(1);
if(x < sx + sw && y < sy + sh)
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y);
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)

8
intern/cycles/kernel/kernel.cpp
View File

@@ -204,16 +204,16 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
/* Path Tracing */
void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y)
void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
{
kernel_path_trace(kg, buffer, rng_state, sample, x, y);
kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
}
/* Tonemapping */
void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y)
void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
{
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y);
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
}
/* Displacement */

8
intern/cycles/kernel/kernel.cu
View File

@@ -26,22 +26,22 @@
#include "kernel_path.h"
#include "kernel_displace.h"
extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh)
extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
if(x < sx + sw && y < sy + sh)
kernel_path_trace(NULL, buffer, rng_state, sample, x, y);
kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
}
extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh)
extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride)
{
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
if(x < sx + sw && y < sy + sh)
kernel_film_tonemap(NULL, rgba, buffer, sample, resolution, x, y);
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)

12
intern/cycles/kernel/kernel.h
View File

@@ -36,13 +36,17 @@ bool kernel_osl_use(KernelGlobals *kg);
void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size);
void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t width, size_t height);
void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y);
void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y);
void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
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);
#ifdef WITH_OPTIMIZED_KERNEL
void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y);
void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y);
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,
int sample, int resolution, int x, int y, int offset, int stride);
void kernel_cpu_optimized_displace(KernelGlobals *kg, uint4 *input, float3 *offset, int i);
#endif

8
intern/cycles/kernel/kernel_camera.h
View File

@@ -74,8 +74,8 @@ __device void camera_sample_perspective(KernelGlobals *kg, float raster_x, float
ray->dP.dx = make_float3(0.0f, 0.0f, 0.0f);
ray->dP.dy = make_float3(0.0f, 0.0f, 0.0f);
ray->dD.dx = normalize(Ddiff + kernel_data.cam.dx) - normalize(Ddiff);
ray->dD.dy = normalize(Ddiff + kernel_data.cam.dy) - normalize(Ddiff);
ray->dD.dx = normalize(Ddiff + float4_to_float3(kernel_data.cam.dx)) - normalize(Ddiff);
ray->dD.dy = normalize(Ddiff + float4_to_float3(kernel_data.cam.dy)) - normalize(Ddiff);
#endif
#ifdef __CAMERA_CLIPPING__
@@ -107,8 +107,8 @@ __device void camera_sample_orthographic(KernelGlobals *kg, float raster_x, floa
#ifdef __RAY_DIFFERENTIALS__
/* ray differential */
ray->dP.dx = kernel_data.cam.dx;
ray->dP.dy = kernel_data.cam.dy;
ray->dP.dx = float4_to_float3(kernel_data.cam.dx);
ray->dP.dy = float4_to_float3(kernel_data.cam.dy);
ray->dD.dx = make_float3(0.0f, 0.0f, 0.0f);
ray->dD.dy = make_float3(0.0f, 0.0f, 0.0f);

19
intern/cycles/kernel/kernel_compat_opencl.h
View File

@@ -25,12 +25,21 @@
/* no namespaces in opencl */
#define CCL_NAMESPACE_BEGIN
#define CCL_NAMESPACE_END
#define WITH_OPENCL
#ifdef __CL_NO_FLOAT3__
#define float3 float4
#endif
#ifdef __CL_NOINLINE__
#define __noinline __attribute__((noinline))
#else
#define __noinline
#endif
/* in opencl all functions are device functions, so leave this empty */
#define __device
#define __device_inline
#define __device_noinline
#define __device_inline __device
#define __device_noinline __device __noinline
/* no assert in opencl */
#define kernel_assert(cond)
@@ -68,7 +77,11 @@ __device float kernel_tex_interp_(__global float *data, int width, float x)
#endif
#define make_float2(x, y) ((float2)(x, y))
#ifdef __CL_NO_FLOAT3__
#define make_float3(x, y, z) ((float4)(x, y, z, 0.0))
#else
#define make_float3(x, y, z) ((float3)(x, y, z))
#endif
#define make_float4(x, y, z, w) ((float4)(x, y, z, w))
#define make_int2(x, y) ((int2)(x, y))
#define make_int3(x, y, z) ((int3)(x, y, z))

5
intern/cycles/kernel/kernel_film.h
View File

@@ -48,10 +48,9 @@ __device uchar4 film_float_to_byte(float4 color)
return result;
}
__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y)
__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
{
int w = kernel_data.cam.width;
int index = x + y*w;
int index = offset + x + y*stride;
float4 irradiance = buffer[index];
float4 float_result = film_map(kg, irradiance, sample);

8
intern/cycles/kernel/kernel_optimized.cpp
View File

@@ -35,16 +35,16 @@ CCL_NAMESPACE_BEGIN
/* Path Tracing */
void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y)
void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
{
kernel_path_trace(kg, buffer, rng_state, sample, x, y);
kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
}
/* Tonemapping */
void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y)
void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
{
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y);
kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
}
/* Displacement */

14
intern/cycles/kernel/kernel_path.h
View File

@@ -34,7 +34,7 @@
CCL_NAMESPACE_BEGIN
#ifdef __MODIFY_TP__
__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int sample)
__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int offset, int stride, int sample)
{
/* modify throughput to influence path termination probability, to avoid
darker regions receiving fewer samples than lighter regions. also RGB
@@ -45,7 +45,7 @@ __device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global f
const float minL = 0.1f;
if(sample >= minsample) {
float3 L = buffer[x + y*kernel_data.cam.width];
float3 L = buffer[offset + x + y*stride];
float3 Lmin = make_float3(minL, minL, minL);
float correct = (float)(sample+1)/(float)sample;
@@ -379,7 +379,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent);
}
__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y)
__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride)
{
/* initialize random numbers */
RNG rng;
@@ -387,7 +387,7 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
float filter_u;
float filter_v;
path_rng_init(kg, rng_state, sample, &rng, x, y, &filter_u, &filter_v);
path_rng_init(kg, rng_state, sample, &rng, x, y, offset, stride, &filter_u, &filter_v);
/* sample camera ray */
Ray ray;
@@ -399,7 +399,7 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
/* integrate */
#ifdef __MODIFY_TP__
float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, sample);
float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, offset, stride, sample);
float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput)/throughput;
#else
float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
@@ -407,14 +407,14 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
#endif
/* accumulate result in output buffer */
int index = x + y*kernel_data.cam.width;
int index = offset + x + y*stride;
if(sample == 0)
buffer[index] = L;
else
buffer[index] += L;
path_rng_end(kg, rng_state, rng, x, y);
path_rng_end(kg, rng_state, rng, x, y, offset, stride);
}
CCL_NAMESPACE_END

14
intern/cycles/kernel/kernel_random.h
View File

@@ -123,7 +123,7 @@ __device_inline float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dime
#endif
}
__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
{
#ifdef __SOBOL_FULL_SCREEN__
uint px, py;
@@ -138,7 +138,7 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state,
*fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x;
*fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y;
#else
*rng = rng_state[x + y*kernel_data.cam.width];
*rng = rng_state[offset + x + y*stride];
*rng ^= kernel_data.integrator.seed;
@@ -147,7 +147,7 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state,
#endif
}
__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y)
__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
{
/* nothing to do */
}
@@ -163,10 +163,10 @@ __device float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dimension)
return (float)*rng * (1.0f/(float)0xFFFFFFFF);
}
__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
{
/* load state */
*rng = rng_state[x + y*kernel_data.cam.width];
*rng = rng_state[offset + x + y*stride];
*rng ^= kernel_data.integrator.seed;
@@ -174,10 +174,10 @@ __device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sam
*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
}
__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y)
__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
{
/* store state for next sample */
rng_state[x + y*kernel_data.cam.width] = rng;
rng_state[offset + x + y*stride] = rng;
}
#endif

25
intern/cycles/kernel/kernel_types.h
View File

@@ -295,29 +295,24 @@ typedef struct ShaderData {
#endif
} ShaderData;
/* Constrant Kernel Data */
/* Constrant Kernel Data
*
* These structs are passed from CPU to various devices, and the struct layout
* must match exactly. Structs are padded to ensure 16 byte alignment, and we
* do not use float3 because its size may not be the same on all devices. */
typedef struct KernelCamera {
/* type */
int ortho;
int pad;
/* size */
int width, height;
int pad1, pad2, pad3;
/* matrices */
Transform cameratoworld;
Transform rastertocamera;
/* differentials */
float3 dx;
#ifndef WITH_OPENCL
float pad1;
#endif
float3 dy;
#ifndef WITH_OPENCL
float pad2;
#endif
float4 dx;
float4 dy;
/* depth of field */
float aperturesize;
@@ -358,10 +353,6 @@ typedef struct KernelBackground {
typedef struct KernelSunSky {
/* sun direction in spherical and cartesian */
float theta, phi, pad3, pad4;
float3 dir;
#ifndef WITH_OPENCL
float pad;
#endif
/* perez function parameters */
float zenith_Y, zenith_x, zenith_y, pad2;

42
intern/cycles/render/buffers.cpp
View File

@@ -36,8 +36,6 @@ CCL_NAMESPACE_BEGIN
RenderBuffers::RenderBuffers(Device *device_)
{
device = device_;
width = 0;
height = 0;
}
RenderBuffers::~RenderBuffers()
@@ -58,24 +56,23 @@ void RenderBuffers::device_free()
}
}
void RenderBuffers::reset(Device *device, int width_, int height_)
void RenderBuffers::reset(Device *device, BufferParams& params_)
{
width = width_;
height = height_;
params = params_;
/* free existing buffers */
device_free();
/* allocate buffer */
buffer.resize(width, height);
buffer.resize(params.width, params.height);
device->mem_alloc(buffer, MEM_READ_WRITE);
device->mem_zero(buffer);
/* allocate rng state */
rng_state.resize(width, height);
rng_state.resize(params.width, params.height);
uint *init_state = rng_state.resize(width, height);
int x, y;
uint *init_state = rng_state.resize(params.width, params.height);
int x, y, width = params.width, height = params.height;
for(x=0; x<width; x++)
for(y=0; y<height; y++)
@@ -92,11 +89,11 @@ float4 *RenderBuffers::copy_from_device(float exposure, int sample)
device->mem_copy_from(buffer, 0, buffer.memory_size());
float4 *out = new float4[width*height];
float4 *out = new float4[params.width*params.height];
float4 *in = (float4*)buffer.data_pointer;
float scale = 1.0f/(float)sample;
for(int i = width*height - 1; i >= 0; i--) {
for(int i = params.width*params.height - 1; i >= 0; i--) {
float4 rgba = in[i]*scale;
rgba.x = rgba.x*exposure;
@@ -117,8 +114,6 @@ float4 *RenderBuffers::copy_from_device(float exposure, int sample)
DisplayBuffer::DisplayBuffer(Device *device_)
{
device = device_;
width = 0;
height = 0;
draw_width = 0;
draw_height = 0;
transparent = true; /* todo: determine from background */
@@ -137,28 +132,27 @@ void DisplayBuffer::device_free()
}
}
void DisplayBuffer::reset(Device *device, int width_, int height_)
void DisplayBuffer::reset(Device *device, BufferParams& params_)
{
draw_width = 0;
draw_height = 0;
width = width_;
height = height_;
params = params_;
/* free existing buffers */
device_free();
/* allocate display pixels */
rgba.resize(width, height);
rgba.resize(params.width, params.height);
device->pixels_alloc(rgba);
}
void DisplayBuffer::draw_set(int width_, int height_)
void DisplayBuffer::draw_set(int width, int height)
{
assert(width_ <= width && height_ <= height);
assert(width <= params.width && height <= params.height);
draw_width = width_;
draw_height = height_;
draw_width = width;
draw_height = height;
}
void DisplayBuffer::draw_transparency_grid()
@@ -175,11 +169,11 @@ void DisplayBuffer::draw_transparency_grid()
};
glColor4ub(50, 50, 50, 255);
glRectf(0, 0, width, height);
glRectf(0, 0, params.width, params.height);
glEnable(GL_POLYGON_STIPPLE);
glColor4ub(55, 55, 55, 255);
glPolygonStipple(checker_stipple_sml);
glRectf(0, 0, width, height);
glRectf(0, 0, params.width, params.height);
glDisable(GL_POLYGON_STIPPLE);
}
@@ -189,7 +183,7 @@ void DisplayBuffer::draw(Device *device)
if(transparent)
draw_transparency_grid();
device->draw_pixels(rgba, 0, draw_width, draw_height, width, height, transparent);
device->draw_pixels(rgba, 0, draw_width, draw_height, params.width, params.height, transparent);
}
}

49
intern/cycles/render/buffers.h
View File

@@ -30,12 +30,49 @@ CCL_NAMESPACE_BEGIN
class Device;
struct float4;
/* Buffer Parameters
Size of render buffer and how it fits in the full image (border render). */
class BufferParams {
public:
/* width/height of the physical buffer */
int width;
int height;
/* offset into and width/height of the full buffer */
int full_x;
int full_y;
int full_width;
int full_height;
BufferParams()
{
width = 0;
height = 0;
full_x = 0;
full_y = 0;
full_width = 0;
full_height = 0;
}
bool modified(const BufferParams& params)
{
return !(full_x == params.full_x
&& full_y == params.full_y
&& width == params.width
&& height == params.height
&& full_width == params.full_width
&& full_height == params.full_height);
}
};
/* Render Buffers */
class RenderBuffers {
public:
/* buffer dimensions */
int width, height;
/* buffer parameters */
BufferParams params;
/* float buffer */
device_vector<float4> buffer;
/* random number generator state */
@@ -46,7 +83,7 @@ public:
RenderBuffers(Device *device);
~RenderBuffers();
void reset(Device *device, int width, int height);
void reset(Device *device, BufferParams& params);
float4 *copy_from_device(float exposure, int sample);
protected:
@@ -62,8 +99,8 @@ protected:
class DisplayBuffer {
public:
/* buffer dimensions */
int width, height;
/* buffer parameters */
BufferParams params;
/* dimensions for how much of the buffer is actually ready for display.
with progressive render we can be using only a subset of the buffer.
if these are zero, it means nothing can be drawn yet */
@@ -78,7 +115,7 @@ public:
DisplayBuffer(Device *device);
~DisplayBuffer();
void reset(Device *device, int width, int height);
void reset(Device *device, BufferParams& params);
void write(Device *device, const string& filename);
void draw_set(int width, int height);

11
intern/cycles/render/camera.cpp
View File

@@ -72,8 +72,9 @@ void Camera::update()
if(!need_update)
return;
/* ndc to raster */
Transform screentocamera;
Transform ndctoraster = transform_scale((float)width, (float)height, 1.0f);
Transform ndctoraster = transform_scale(width, height, 1.0f);
/* raster to screen */
Transform screentoraster = ndctoraster *
@@ -148,13 +149,9 @@ void Camera::device_update(Device *device, DeviceScene *dscene)
/* type */
kcam->ortho = ortho;
/* size */
kcam->width = width;
kcam->height = height;
/* store differentials */
kcam->dx = dx;
kcam->dy = dy;
kcam->dx = float3_to_float4(dx);
kcam->dy = float3_to_float4(dy);
/* clipping */
kcam->nearclip = nearclip;

1
intern/cycles/render/nodes.cpp
View File

@@ -273,7 +273,6 @@ static void sky_texture_precompute(KernelSunSky *ksunsky, float3 dir, float turb
ksunsky->theta = theta;
ksunsky->phi = phi;
ksunsky->dir = dir;
float theta2 = theta*theta;
float theta3 = theta*theta*theta;

73
intern/cycles/render/session.cpp
View File

@@ -51,8 +51,6 @@ Session::Session(const SessionParams& params_)
sample = 0;
delayed_reset.do_reset = false;
delayed_reset.w = 0;
delayed_reset.h = 0;
delayed_reset.samples = 0;
display_outdated = false;
@@ -108,7 +106,7 @@ bool Session::ready_to_reset()
/* GPU Session */
void Session::reset_gpu(int w, int h, int samples)
void Session::reset_gpu(BufferParams& buffer_params, int samples)
{
/* block for buffer acces and reset immediately. we can't do this
in the thread, because we need to allocate an OpenGL buffer, and
@@ -119,7 +117,7 @@ void Session::reset_gpu(int w, int h, int samples)
display_outdated = true;
reset_time = time_dt();
reset_(w, h, samples);
reset_(buffer_params, samples);
gpu_need_tonemap = false;
gpu_need_tonemap_cond.notify_all();
@@ -127,7 +125,7 @@ void Session::reset_gpu(int w, int h, int samples)
pause_cond.notify_all();
}
bool Session::draw_gpu(int w, int h)
bool Session::draw_gpu(BufferParams& buffer_params)
{
/* block for buffer access */
thread_scoped_lock display_lock(display->mutex);
@@ -136,7 +134,7 @@ bool Session::draw_gpu(int w, int h)
if(gpu_draw_ready) {
/* then verify the buffers have the expected size, so we don't
draw previous results in a resized window */
if(w == display->width && h == display->height) {
if(!buffer_params.modified(display->params)) {
/* for CUDA we need to do tonemapping still, since we can
only access GL buffers from the main thread */
if(gpu_need_tonemap) {
@@ -261,15 +259,14 @@ void Session::run_gpu()
/* CPU Session */
void Session::reset_cpu(int w, int h, int samples)
void Session::reset_cpu(BufferParams& buffer_params, int samples)
{
thread_scoped_lock reset_lock(delayed_reset.mutex);
display_outdated = true;
reset_time = time_dt();
delayed_reset.w = w;
delayed_reset.h = h;
delayed_reset.params = buffer_params;
delayed_reset.samples = samples;
delayed_reset.do_reset = true;
device->task_cancel();
@@ -277,7 +274,7 @@ void Session::reset_cpu(int w, int h, int samples)
pause_cond.notify_all();
}
bool Session::draw_cpu(int w, int h)
bool Session::draw_cpu(BufferParams& buffer_params)
{
thread_scoped_lock display_lock(display->mutex);
@@ -285,7 +282,7 @@ bool Session::draw_cpu(int w, int h)
if(display->draw_ready()) {
/* then verify the buffers have the expected size, so we don't
draw previous results in a resized window */
if(w == display->width && h == display->height) {
if(!buffer_params.modified(display->params)) {
display->draw(device);
if(display_outdated && (time_dt() - reset_time) > params.text_timeout)
@@ -306,7 +303,7 @@ void Session::run_cpu()
thread_scoped_lock buffers_lock(buffers->mutex);
thread_scoped_lock display_lock(display->mutex);
reset_(delayed_reset.w, delayed_reset.h, delayed_reset.samples);
reset_(delayed_reset.params, delayed_reset.samples);
delayed_reset.do_reset = false;
}
@@ -389,7 +386,7 @@ void Session::run_cpu()
if(delayed_reset.do_reset) {
/* reset rendering if request from main thread */
delayed_reset.do_reset = false;
reset_(delayed_reset.w, delayed_reset.h, delayed_reset.samples);
reset_(delayed_reset.params, delayed_reset.samples);
}
else if(need_tonemap) {
/* tonemap only if we do not reset, we don't we don't
@@ -438,23 +435,23 @@ void Session::run()
progress.set_update();
}
bool Session::draw(int w, int h)
bool Session::draw(BufferParams& buffer_params)
{
if(device_use_gl)
return draw_gpu(w, h);
return draw_gpu(buffer_params);
else
return draw_cpu(w, h);
return draw_cpu(buffer_params);
}
void Session::reset_(int w, int h, int samples)
void Session::reset_(BufferParams& buffer_params, int samples)
{
if(w != buffers->width || h != buffers->height) {
if(buffer_params.modified(buffers->params)) {
gpu_draw_ready = false;
buffers->reset(device, w, h);
display->reset(device, w, h);
buffers->reset(device, buffer_params);
display->reset(device, buffer_params);
}
tile_manager.reset(w, h, samples);
tile_manager.reset(buffer_params, samples);
start_time = time_dt();
preview_time = 0.0;
@@ -462,12 +459,12 @@ void Session::reset_(int w, int h, int samples)
sample = 0;
}
void Session::reset(int w, int h, int samples)
void Session::reset(BufferParams& buffer_params, int samples)
{
if(device_use_gl)
reset_gpu(w, h, samples);
reset_gpu(buffer_params, samples);
else
reset_cpu(w, h, samples);
reset_cpu(buffer_params, samples);
}
void Session::set_samples(int samples)
@@ -514,14 +511,18 @@ void Session::update_scene()
progress.set_status("Updating Scene");
/* update camera if dimensions changed for progressive render */
/* update camera if dimensions changed for progressive render. the camera
knows nothing about progressive or cropped rendering, it just gets the
image dimensions passed in */
Camera *cam = scene->camera;
int w = tile_manager.state.width;
int h = tile_manager.state.height;
float progressive_x = tile_manager.state.width/(float)tile_manager.params.width;
float progressive_y = tile_manager.state.height/(float)tile_manager.params.height;
int width = tile_manager.params.full_width*progressive_x;
int height = tile_manager.params.full_height*progressive_y;
if(cam->width != w || cam->height != h) {
cam->width = w;
cam->height = h;
if(width != cam->width || height != cam->height) {
cam->width = width;
cam->height = height;
cam->tag_update();
}
@@ -573,14 +574,16 @@ void Session::path_trace(Tile& tile)
/* add path trace task */
DeviceTask task(DeviceTask::PATH_TRACE);
task.x = tile.x;
task.y = tile.y;
task.x = tile_manager.state.full_x + tile.x;
task.y = tile_manager.state.full_y + tile.y;
task.w = tile.w;
task.h = tile.h;
task.buffer = buffers->buffer.device_pointer;
task.rng_state = buffers->rng_state.device_pointer;
task.sample = tile_manager.state.sample;
task.resolution = tile_manager.state.resolution;
task.offset = -(tile_manager.state.full_x + tile_manager.state.full_y*tile_manager.state.width);
task.stride = tile_manager.state.width;
device->task_add(task);
}
@@ -590,14 +593,16 @@ void Session::tonemap()
/* add tonemap task */
DeviceTask task(DeviceTask::TONEMAP);
task.x = 0;
task.y = 0;
task.x = tile_manager.state.full_x;
task.y = tile_manager.state.full_y;
task.w = tile_manager.state.width;
task.h = tile_manager.state.height;
task.rgba = display->rgba.device_pointer;
task.buffer = buffers->buffer.device_pointer;
task.sample = tile_manager.state.sample;
task.resolution = tile_manager.state.resolution;
task.offset = -(tile_manager.state.full_x + tile_manager.state.full_y*tile_manager.state.width);
task.stride = tile_manager.state.width;
if(task.w > 0 && task.h > 0) {
device->task_add(task);

18
intern/cycles/render/session.h
View File

@@ -19,6 +19,7 @@
#ifndef __SESSION_H__
#define __SESSION_H__
#include "buffers.h"
#include "device.h"
#include "tile.h"
@@ -27,6 +28,7 @@
CCL_NAMESPACE_BEGIN
class BufferParams;
class Device;
class DeviceScene;
class DisplayBuffer;
@@ -106,11 +108,11 @@ public:
~Session();
void start();
bool draw(int w, int h);
bool draw(BufferParams& params);
void wait();
bool ready_to_reset();
void reset(int w, int h, int samples);
void reset(BufferParams& params, int samples);
void set_samples(int samples);
void set_pause(bool pause);
@@ -118,7 +120,7 @@ protected:
struct DelayedReset {
thread_mutex mutex;
bool do_reset;
int w, h;
BufferParams params;
int samples;
} delayed_reset;
@@ -129,15 +131,15 @@ protected:
void tonemap();
void path_trace(Tile& tile);
void reset_(int w, int h, int samples);
void reset_(BufferParams& params, int samples);
void run_cpu();
bool draw_cpu(int w, int h);
void reset_cpu(int w, int h, int samples);
bool draw_cpu(BufferParams& params);
void reset_cpu(BufferParams& params, int samples);
void run_gpu();
bool draw_gpu(int w, int h);
void reset_gpu(int w, int h, int samples);
bool draw_gpu(BufferParams& params);
void reset_gpu(BufferParams& params, int samples);
TileManager tile_manager;
bool device_use_gl;

18
intern/cycles/render/tile.cpp
View File

@@ -28,21 +28,21 @@ TileManager::TileManager(bool progressive_, int samples_, int tile_size_, int mi
tile_size = tile_size_;
min_size = min_size_;
reset(0, 0, 0);
BufferParams buffer_params;
reset(buffer_params, 0);
}
TileManager::~TileManager()
{
}
void TileManager::reset(int width_, int height_, int samples_)
void TileManager::reset(BufferParams& params_, int samples_)
{
full_width = width_;
full_height = height_;
params = params_;
start_resolution = 1;
int w = width_, h = height_;
int w = params.width, h = params.height;
if(min_size != INT_MAX) {
while(w*h > min_size*min_size) {
@@ -55,6 +55,8 @@ void TileManager::reset(int width_, int height_, int samples_)
samples = samples_;
state.full_x = 0;
state.full_y = 0;
state.width = 0;
state.height = 0;
state.sample = -1;
@@ -70,8 +72,8 @@ void TileManager::set_samples(int samples_)
void TileManager::set_tiles()
{
int resolution = state.resolution;
int image_w = max(1, full_width/resolution);
int image_h = max(1, full_height/resolution);
int image_w = max(1, params.width/resolution);
int image_h = max(1, params.height/resolution);
int tile_w = (image_w + tile_size - 1)/tile_size;
int tile_h = (image_h + tile_size - 1)/tile_size;
int sub_w = image_w/tile_w;
@@ -90,6 +92,8 @@ void TileManager::set_tiles()
}
}
state.full_x = params.full_x/resolution;
state.full_y = params.full_y/resolution;
state.width = image_w;
state.height = image_h;
}

8
intern/cycles/render/tile.h
View File

@@ -21,6 +21,7 @@
#include <limits.h>
#include "buffers.h"
#include "util_list.h"
CCL_NAMESPACE_BEGIN
@@ -39,7 +40,10 @@ public:
class TileManager {
public:
BufferParams params;
struct State {
int full_x;
int full_y;
int width;
int height;
int sample;
@@ -50,7 +54,7 @@ public:
TileManager(bool progressive, int samples, int tile_size, int min_size);
~TileManager();
void reset(int width, int height, int samples);
void reset(BufferParams& params, int samples);
void set_samples(int samples);
bool next();
bool done();
@@ -63,8 +67,6 @@ protected:
int tile_size;
int min_size;
int full_width;
int full_height;
int start_resolution;
};

5
intern/cycles/util/util_math.h
View File

@@ -536,6 +536,11 @@ __device_inline float3 float4_to_float3(const float4 a)
return make_float3(a.x, a.y, a.z);
}
__device_inline float4 float3_to_float4(const float3 a)
{
return make_float4(a.x, a.y, a.z, 1.0f);
}
#ifndef __KERNEL_GPU__
__device_inline void print_float3(const char *label, const float3& a)

6
source/blender/blenkernel/BKE_modifier.h
View File

@@ -66,6 +66,11 @@ typedef enum {
* unless it's a mesh and can be exploded -> curve can also emit particles
*/
eModifierTypeType_DeformOrConstruct,
/* Like eModifierTypeType_Nonconstructive, but does not affect the geometry
* of the object, rather some of its CustomData layers.
* E.g. UVProject and WeightVG modifiers. */
eModifierTypeType_NonGeometrical,
} ModifierTypeType;
typedef enum {
@@ -312,6 +317,7 @@ int modifier_supportsMapping(struct ModifierData *md);
int modifier_couldBeCage(struct Scene *scene, struct ModifierData *md);
int modifier_isCorrectableDeformed(struct ModifierData *md);
int modifier_sameTopology(ModifierData *md);
int modifier_nonGeometrical(ModifierData *md);
int modifier_isEnabled(struct Scene *scene, struct ModifierData *md, int required_mode);
void modifier_setError(struct ModifierData *md, const char *format, ...)
#ifdef __GNUC__

9
source/blender/blenkernel/intern/modifier.c
View File

@@ -239,7 +239,14 @@ int modifier_couldBeCage(struct Scene *scene, ModifierData *md)
int modifier_sameTopology(ModifierData *md)
{
ModifierTypeInfo *mti = modifierType_getInfo(md->type);
return ( mti->type == eModifierTypeType_OnlyDeform || mti->type == eModifierTypeType_Nonconstructive);
return ELEM3(mti->type, eModifierTypeType_OnlyDeform, eModifierTypeType_Nonconstructive,
eModifierTypeType_NonGeometrical);
}
int modifier_nonGeometrical(ModifierData *md)
{
ModifierTypeInfo *mti = modifierType_getInfo(md->type);
return (mti->type == eModifierTypeType_NonGeometrical);
}
void modifier_setError(ModifierData *md, const char *format, ...)

4
source/blender/editors/interface/interface_templates.c
View File

@@ -845,7 +845,7 @@ static uiLayout *draw_modifier(uiLayout *layout, Scene *scene, Object *ob, Modif
uiLayoutSetOperatorContext(row, WM_OP_INVOKE_DEFAULT);
uiItemEnumO(row, "OBJECT_OT_modifier_apply", IFACE_("Apply"), 0, "apply_as", MODIFIER_APPLY_DATA);
if (modifier_sameTopology(md))
if (modifier_sameTopology(md) && !modifier_nonGeometrical(md))
uiItemEnumO(row, "OBJECT_OT_modifier_apply", IFACE_("Apply as Shape"), 0, "apply_as", MODIFIER_APPLY_SHAPE);
}
@@ -853,7 +853,7 @@ static uiLayout *draw_modifier(uiLayout *layout, Scene *scene, Object *ob, Modif
uiBlockSetButLock(block, ob && ob->id.lib, ERROR_LIBDATA_MESSAGE);
if (!ELEM5(md->type, eModifierType_Fluidsim, eModifierType_Softbody, eModifierType_ParticleSystem, eModifierType_Cloth, eModifierType_Smoke))
uiItemO(row, TIP_("Copy"), ICON_NONE, "OBJECT_OT_modifier_copy");
uiItemO(row, IFACE_("Copy"), ICON_NONE, "OBJECT_OT_modifier_copy");
}
/* result is the layout block inside the box, that we return so that modifier settings can be drawn */

4
source/blender/editors/object/object_modifier.c
View File

@@ -464,7 +464,7 @@ static int modifier_apply_shape(ReportList *reports, Scene *scene, Object *ob, M
Key *key=me->key;
KeyBlock *kb;
if(!modifier_sameTopology(md)) {
if(!modifier_sameTopology(md) || mti->type == eModifierTypeType_NonGeometrical) {
BKE_report(reports, RPT_ERROR, "Only deforming modifiers can be applied to Shapes");
return 0;
}
@@ -512,7 +512,7 @@ static int modifier_apply_obdata(ReportList *reports, Scene *scene, Object *ob,
Mesh *me = ob->data;
MultiresModifierData *mmd= find_multires_modifier_before(scene, md);
if( me->key) {
if(me->key && mti->type != eModifierTypeType_NonGeometrical) {
BKE_report(reports, RPT_ERROR, "Modifier cannot be applied to Mesh with Shape Keys");
return 0;
}

2
source/blender/modifiers/intern/MOD_uvproject.c
View File

@@ -408,7 +408,7 @@ ModifierTypeInfo modifierType_UVProject = {
/* name */ "UVProject",
/* structName */ "UVProjectModifierData",
/* structSize */ sizeof(UVProjectModifierData),
/* type */ eModifierTypeType_Nonconstructive,
/* type */ eModifierTypeType_NonGeometrical,
/* flags */ eModifierTypeFlag_AcceptsMesh
| eModifierTypeFlag_SupportsMapping
| eModifierTypeFlag_SupportsEditmode

2
source/blender/modifiers/intern/MOD_weightvgedit.c
View File

@@ -254,7 +254,7 @@ ModifierTypeInfo modifierType_WeightVGEdit = {
/* name */ "VertexWeightEdit",
/* structName */ "WeightVGEditModifierData",
/* structSize */ sizeof(WeightVGEditModifierData),
/* type */ eModifierTypeType_Nonconstructive,
/* type */ eModifierTypeType_NonGeometrical,
/* flags */ eModifierTypeFlag_AcceptsMesh
/* |eModifierTypeFlag_SupportsMapping*/
|eModifierTypeFlag_SupportsEditmode,

2
source/blender/modifiers/intern/MOD_weightvgmix.c
View File

@@ -386,7 +386,7 @@ ModifierTypeInfo modifierType_WeightVGMix = {
/* name */ "VertexWeightMix",
/* structName */ "WeightVGMixModifierData",
/* structSize */ sizeof(WeightVGMixModifierData),
/* type */ eModifierTypeType_Nonconstructive,
/* type */ eModifierTypeType_NonGeometrical,
/* flags */ eModifierTypeFlag_AcceptsMesh
/* |eModifierTypeFlag_SupportsMapping*/
|eModifierTypeFlag_SupportsEditmode,

2
source/blender/modifiers/intern/MOD_weightvgproximity.c
View File

@@ -520,7 +520,7 @@ ModifierTypeInfo modifierType_WeightVGProximity = {
/* name */ "VertexWeightProximity",
/* structName */ "WeightVGProximityModifierData",
/* structSize */ sizeof(WeightVGProximityModifierData),
/* type */ eModifierTypeType_Nonconstructive,
/* type */ eModifierTypeType_NonGeometrical,
/* flags */ eModifierTypeFlag_AcceptsMesh
/* |eModifierTypeFlag_SupportsMapping*/
|eModifierTypeFlag_SupportsEditmode,

16
source/blender/python/mathutils/mathutils_Vector.c
View File

@@ -1486,11 +1486,11 @@ static PyObject *Vector_isub(PyObject *v1, PyObject *v2)
* note: vector/matrix multiplication IS NOT COMMUTATIVE!!!!
* note: assume read callbacks have been done first.
*/
int column_vector_multiplication(float rvec[MAX_DIMENSIONS], VectorObject* vec, MatrixObject * mat)
int column_vector_multiplication(float r_vec[MAX_DIMENSIONS], VectorObject* vec, MatrixObject * mat)
{
float vec_cpy[MAX_DIMENSIONS];
double dot = 0.0f;
int x, y, z = 0;
int row, col, z = 0;
if (mat->num_col != vec->size) {
if (mat->num_col == 4 && vec->size == 3) {
@@ -1507,13 +1507,13 @@ int column_vector_multiplication(float rvec[MAX_DIMENSIONS], VectorObject* vec,
memcpy(vec_cpy, vec->vec, vec->size * sizeof(float));
rvec[3] = 1.0f;
r_vec[3] = 1.0f;
for (x = 0; x < mat->num_row; x++) {
for (y = 0; y < mat->num_col; y++) {
dot += (double)(MATRIX_ITEM(mat, y, x) * vec_cpy[y]);
for (row = 0; row < mat->num_row; row++) {
for (col = 0; col < mat->num_col; col++) {
dot += (double)(MATRIX_ITEM(mat, row, col) * vec_cpy[col]);
}
rvec[z++] = (float)dot;
r_vec[z++] = (float)dot;
dot = 0.0f;
}
@@ -2634,7 +2634,7 @@ static int row_vector_multiplication(float rvec[MAX_DIMENSIONS], VectorObject *v
//muliplication
for (x = 0; x < mat->num_col; x++) {
for (y = 0; y < mat->num_row; y++) {
dot += MATRIX_ITEM(mat, x, y) * vec_cpy[y];
dot += MATRIX_ITEM(mat, y, x) * vec_cpy[y];
}
rvec[z++] = (float)dot;
dot = 0.0f;

5
source/blender/render/intern/source/external_engine.c
View File

@@ -190,6 +190,11 @@ void RE_engine_end_result(RenderEngine *engine, RenderResult *result)
if(!result)
return;
result->tilerect.xmin += re->disprect.xmin;
result->tilerect.xmax += re->disprect.xmin;
result->tilerect.ymin += re->disprect.ymin;
result->tilerect.ymax += re->disprect.ymin;
/* merge. on break, don't merge in result for preview renders, looks nicer */
if(!(re->test_break(re->tbh) && (re->r.scemode & R_PREVIEWBUTS)))