forked from bartvdbraak/blender
Code refactor: move more memory allocation logic into device API.
* Remove tex_* and pixels_* functions, replace by mem_*. * Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices. * No longer create device_memory and call mem_* directly, always go through device_only_memory, device_vector and device_pixels.
This commit is contained in:
parent
aa8b4c5d81
commit
070a668d04
@ -85,28 +85,12 @@ Device::~Device()
|
||||
}
|
||||
}
|
||||
|
||||
void Device::pixels_alloc(device_memory& mem)
|
||||
{
|
||||
mem_alloc(mem);
|
||||
}
|
||||
|
||||
void Device::pixels_copy_from(device_memory& mem, int y, int w, int h)
|
||||
{
|
||||
if(mem.data_type == TYPE_HALF)
|
||||
mem_copy_from(mem, y, w, h, sizeof(half4));
|
||||
else
|
||||
mem_copy_from(mem, y, w, h, sizeof(uchar4));
|
||||
}
|
||||
|
||||
void Device::pixels_free(device_memory& mem)
|
||||
{
|
||||
mem_free(mem);
|
||||
}
|
||||
|
||||
void Device::draw_pixels(device_memory& rgba, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
|
||||
const DeviceDrawParams &draw_params)
|
||||
{
|
||||
pixels_copy_from(rgba, y, w, h);
|
||||
assert(mem.type == MEM_PIXELS);
|
||||
|
||||
mem_copy_from(rgba, y, w, h, rgba.memory_elements_size(1));
|
||||
|
||||
if(transparent) {
|
||||
glEnable(GL_BLEND);
|
||||
|
@ -281,28 +281,12 @@ public:
|
||||
/* statistics */
|
||||
Stats &stats;
|
||||
|
||||
/* regular memory */
|
||||
virtual void mem_alloc(device_memory& mem) = 0;
|
||||
virtual void mem_copy_to(device_memory& mem) = 0;
|
||||
virtual void mem_copy_from(device_memory& mem,
|
||||
int y, int w, int h, int elem) = 0;
|
||||
virtual void mem_zero(device_memory& mem) = 0;
|
||||
virtual void mem_free(device_memory& mem) = 0;
|
||||
|
||||
/* memory alignment */
|
||||
virtual int mem_address_alignment() { return 16; }
|
||||
|
||||
/* constant memory */
|
||||
virtual void const_copy_to(const char *name, void *host, size_t size) = 0;
|
||||
|
||||
/* texture memory */
|
||||
virtual void tex_alloc(device_memory& /*mem*/) {};
|
||||
virtual void tex_free(device_memory& /*mem*/) {};
|
||||
|
||||
/* pixel memory */
|
||||
virtual void pixels_alloc(device_memory& mem);
|
||||
virtual void pixels_copy_from(device_memory& mem, int y, int w, int h);
|
||||
virtual void pixels_free(device_memory& mem);
|
||||
|
||||
/* open shading language, only for CPU device */
|
||||
virtual void *osl_memory() { return NULL; }
|
||||
|
||||
@ -349,6 +333,20 @@ public:
|
||||
static void tag_update();
|
||||
|
||||
static void free_memory();
|
||||
|
||||
protected:
|
||||
/* Memory allocation, only accessed through device_memory. */
|
||||
friend class MultiDevice;
|
||||
friend class DeviceServer;
|
||||
friend class device_memory;
|
||||
|
||||
virtual void mem_alloc(device_memory& mem) = 0;
|
||||
virtual void mem_copy_to(device_memory& mem) = 0;
|
||||
virtual void mem_copy_from(device_memory& mem,
|
||||
int y, int w, int h, int elem) = 0;
|
||||
virtual void mem_zero(device_memory& mem) = 0;
|
||||
virtual void mem_free(device_memory& mem) = 0;
|
||||
|
||||
private:
|
||||
/* Indicted whether device types and devices lists were initialized. */
|
||||
static bool need_types_update, need_devices_update;
|
||||
|
@ -209,7 +209,7 @@ public:
|
||||
|
||||
CPUDevice(DeviceInfo& info_, Stats &stats_, bool background_)
|
||||
: Device(info_, stats_, background_),
|
||||
texture_info(this, "__texture_info"),
|
||||
texture_info(this, "__texture_info", MEM_TEXTURE),
|
||||
#define REGISTER_KERNEL(name) name ## _kernel(KERNEL_FUNCTIONS(name))
|
||||
REGISTER_KERNEL(path_trace),
|
||||
REGISTER_KERNEL(convert_to_half_float),
|
||||
@ -269,7 +269,7 @@ public:
|
||||
~CPUDevice()
|
||||
{
|
||||
task_pool.stop();
|
||||
tex_free(texture_info);
|
||||
texture_info.free();
|
||||
}
|
||||
|
||||
virtual bool show_samples() const
|
||||
@ -280,33 +280,50 @@ public:
|
||||
void load_texture_info()
|
||||
{
|
||||
if(need_texture_info) {
|
||||
tex_free(texture_info);
|
||||
tex_alloc(texture_info);
|
||||
texture_info.copy_to_device();
|
||||
need_texture_info = false;
|
||||
}
|
||||
}
|
||||
|
||||
void mem_alloc(device_memory& mem)
|
||||
{
|
||||
if(mem.name) {
|
||||
VLOG(1) << "Buffer allocate: " << mem.name << ", "
|
||||
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
|
||||
<< string_human_readable_size(mem.memory_size()) << ")";
|
||||
if(mem.type == MEM_TEXTURE) {
|
||||
assert(!"mem_alloc not supported for textures.");
|
||||
}
|
||||
else {
|
||||
if(mem.name) {
|
||||
VLOG(1) << "Buffer allocate: " << mem.name << ", "
|
||||
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
|
||||
<< string_human_readable_size(mem.memory_size()) << ")";
|
||||
}
|
||||
|
||||
mem.device_pointer = mem.data_pointer;
|
||||
mem.device_pointer = mem.data_pointer;
|
||||
|
||||
if(!mem.device_pointer) {
|
||||
mem.device_pointer = (device_ptr)malloc(mem.memory_size());
|
||||
if(!mem.device_pointer) {
|
||||
mem.device_pointer = (device_ptr)malloc(mem.memory_size());
|
||||
}
|
||||
|
||||
mem.device_size = mem.memory_size();
|
||||
stats.mem_alloc(mem.device_size);
|
||||
}
|
||||
|
||||
mem.device_size = mem.memory_size();
|
||||
stats.mem_alloc(mem.device_size);
|
||||
}
|
||||
|
||||
void mem_copy_to(device_memory& /*mem*/)
|
||||
void mem_copy_to(device_memory& mem)
|
||||
{
|
||||
/* no-op */
|
||||
if(mem.type == MEM_TEXTURE) {
|
||||
tex_free(mem);
|
||||
tex_alloc(mem);
|
||||
}
|
||||
else if(mem.type == MEM_PIXELS) {
|
||||
assert(!"mem_copy_to not supported for pixels.");
|
||||
}
|
||||
else {
|
||||
if(!mem.device_pointer) {
|
||||
mem_alloc(mem);
|
||||
}
|
||||
|
||||
/* copy is no-op */
|
||||
}
|
||||
}
|
||||
|
||||
void mem_copy_from(device_memory& /*mem*/,
|
||||
@ -318,12 +335,21 @@ public:
|
||||
|
||||
void mem_zero(device_memory& mem)
|
||||
{
|
||||
memset((void*)mem.device_pointer, 0, mem.memory_size());
|
||||
if(!mem.device_pointer) {
|
||||
mem_alloc(mem);
|
||||
}
|
||||
|
||||
if(mem.device_pointer) {
|
||||
memset((void*)mem.device_pointer, 0, mem.memory_size());
|
||||
}
|
||||
}
|
||||
|
||||
void mem_free(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
if(mem.type == MEM_TEXTURE) {
|
||||
tex_free(mem);
|
||||
}
|
||||
else if(mem.device_pointer) {
|
||||
if(!mem.data_pointer) {
|
||||
free((void*)mem.device_pointer);
|
||||
}
|
||||
@ -354,7 +380,7 @@ public:
|
||||
kernel_tex_copy(&kernel_globals,
|
||||
mem.name,
|
||||
mem.data_pointer,
|
||||
mem.data_width);
|
||||
mem.data_size);
|
||||
}
|
||||
else {
|
||||
/* Image Texture. */
|
||||
@ -431,13 +457,13 @@ public:
|
||||
|
||||
bool denoising_set_tiles(device_ptr *buffers, DenoisingTask *task)
|
||||
{
|
||||
mem_alloc(task->tiles_mem);
|
||||
|
||||
TilesInfo *tiles = (TilesInfo*) task->tiles_mem.data_pointer;
|
||||
for(int i = 0; i < 9; i++) {
|
||||
tiles->buffers[i] = buffers[i];
|
||||
}
|
||||
|
||||
task->tiles_mem.copy_to_device();
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
@ -723,8 +749,7 @@ public:
|
||||
|
||||
/* allocate buffer for kernel globals */
|
||||
device_only_memory<KernelGlobals> kgbuffer(this, "kernel_globals");
|
||||
kgbuffer.resize(1);
|
||||
mem_alloc(kgbuffer);
|
||||
kgbuffer.alloc_to_device(1);
|
||||
|
||||
KernelGlobals *kg = new ((void*) kgbuffer.device_pointer) KernelGlobals(thread_kernel_globals_init());
|
||||
|
||||
@ -734,8 +759,7 @@ public:
|
||||
requested_features.max_closure = MAX_CLOSURE;
|
||||
if(!split_kernel->load_kernels(requested_features)) {
|
||||
thread_kernel_globals_free((KernelGlobals*)kgbuffer.device_pointer);
|
||||
mem_free(kgbuffer);
|
||||
|
||||
kgbuffer.free();
|
||||
delete split_kernel;
|
||||
return;
|
||||
}
|
||||
@ -766,7 +790,7 @@ public:
|
||||
|
||||
thread_kernel_globals_free((KernelGlobals*)kgbuffer.device_pointer);
|
||||
kg->~KernelGlobals();
|
||||
mem_free(kgbuffer);
|
||||
kgbuffer.free();
|
||||
delete split_kernel;
|
||||
}
|
||||
|
||||
|
@ -218,7 +218,7 @@ public:
|
||||
|
||||
CUDADevice(DeviceInfo& info, Stats &stats, bool background_)
|
||||
: Device(info, stats, background_),
|
||||
texture_info(this, "__texture_info")
|
||||
texture_info(this, "__texture_info", MEM_TEXTURE)
|
||||
{
|
||||
first_error = true;
|
||||
background = background_;
|
||||
@ -275,7 +275,7 @@ public:
|
||||
delete split_kernel;
|
||||
|
||||
if(info.has_bindless_textures) {
|
||||
tex_free(texture_info);
|
||||
texture_info.free();
|
||||
}
|
||||
|
||||
cuda_assert(cuCtxDestroy(cuContext));
|
||||
@ -548,20 +548,19 @@ public:
|
||||
void load_texture_info()
|
||||
{
|
||||
if(info.has_bindless_textures && need_texture_info) {
|
||||
tex_free(texture_info);
|
||||
tex_alloc(texture_info);
|
||||
texture_info.copy_to_device();
|
||||
need_texture_info = false;
|
||||
}
|
||||
}
|
||||
|
||||
void mem_alloc(device_memory& mem)
|
||||
void generic_alloc(device_memory& mem)
|
||||
{
|
||||
CUDAContextScope scope(this);
|
||||
|
||||
if(mem.name) {
|
||||
VLOG(1) << "Buffer allocate: " << mem.name << ", "
|
||||
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
|
||||
<< string_human_readable_size(mem.memory_size()) << ")";
|
||||
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
|
||||
<< string_human_readable_size(mem.memory_size()) << ")";
|
||||
}
|
||||
|
||||
CUdeviceptr device_pointer;
|
||||
@ -572,31 +571,88 @@ public:
|
||||
stats.mem_alloc(size);
|
||||
}
|
||||
|
||||
void generic_copy_to(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
CUDAContextScope scope(this);
|
||||
cuda_assert(cuMemcpyHtoD(cuda_device_ptr(mem.device_pointer), (void*)mem.data_pointer, mem.memory_size()));
|
||||
}
|
||||
}
|
||||
|
||||
void generic_free(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
CUDAContextScope scope(this);
|
||||
|
||||
cuda_assert(cuMemFree(cuda_device_ptr(mem.device_pointer)));
|
||||
|
||||
mem.device_pointer = 0;
|
||||
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_size = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void mem_alloc(device_memory& mem)
|
||||
{
|
||||
if(mem.type == MEM_PIXELS && !background) {
|
||||
pixels_alloc(mem);
|
||||
}
|
||||
else if(mem.type == MEM_TEXTURE) {
|
||||
assert(!"mem_alloc not supported for textures.");
|
||||
}
|
||||
else {
|
||||
generic_alloc(mem);
|
||||
}
|
||||
}
|
||||
|
||||
void mem_copy_to(device_memory& mem)
|
||||
{
|
||||
CUDAContextScope scope(this);
|
||||
if(mem.type == MEM_PIXELS) {
|
||||
assert(!"mem_copy_to not supported for pixels.");
|
||||
}
|
||||
else if(mem.type == MEM_TEXTURE) {
|
||||
tex_free(mem);
|
||||
tex_alloc(mem);
|
||||
}
|
||||
else {
|
||||
if(!mem.device_pointer) {
|
||||
generic_alloc(mem);
|
||||
}
|
||||
|
||||
if(mem.device_pointer)
|
||||
cuda_assert(cuMemcpyHtoD(cuda_device_ptr(mem.device_pointer), (void*)mem.data_pointer, mem.memory_size()));
|
||||
generic_copy_to(mem);
|
||||
}
|
||||
}
|
||||
|
||||
void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
|
||||
{
|
||||
CUDAContextScope scope(this);
|
||||
size_t offset = elem*y*w;
|
||||
size_t size = elem*w*h;
|
||||
|
||||
if(mem.device_pointer) {
|
||||
cuda_assert(cuMemcpyDtoH((uchar*)mem.data_pointer + offset,
|
||||
(CUdeviceptr)(mem.device_pointer + offset), size));
|
||||
if(mem.type == MEM_PIXELS && !background) {
|
||||
pixels_copy_from(mem, y, w, h);
|
||||
}
|
||||
else if(mem.type == MEM_TEXTURE) {
|
||||
assert(!"mem_copy_from not supported for textures.");
|
||||
}
|
||||
else {
|
||||
memset((char*)mem.data_pointer + offset, 0, size);
|
||||
CUDAContextScope scope(this);
|
||||
size_t offset = elem*y*w;
|
||||
size_t size = elem*w*h;
|
||||
|
||||
if(mem.device_pointer) {
|
||||
cuda_assert(cuMemcpyDtoH((uchar*)mem.data_pointer + offset,
|
||||
(CUdeviceptr)(mem.device_pointer + offset), size));
|
||||
}
|
||||
else {
|
||||
memset((char*)mem.data_pointer + offset, 0, size);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void mem_zero(device_memory& mem)
|
||||
{
|
||||
if(!mem.device_pointer) {
|
||||
mem_alloc(mem);
|
||||
}
|
||||
|
||||
if(mem.data_pointer) {
|
||||
memset((void*)mem.data_pointer, 0, mem.memory_size());
|
||||
}
|
||||
@ -609,14 +665,14 @@ public:
|
||||
|
||||
void mem_free(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
CUDAContextScope scope(this);
|
||||
cuda_assert(cuMemFree(cuda_device_ptr(mem.device_pointer)));
|
||||
|
||||
mem.device_pointer = 0;
|
||||
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_size = 0;
|
||||
if(mem.type == MEM_PIXELS && !background) {
|
||||
pixels_free(mem);
|
||||
}
|
||||
else if(mem.type == MEM_TEXTURE) {
|
||||
tex_free(mem);
|
||||
}
|
||||
else {
|
||||
generic_free(mem);
|
||||
}
|
||||
}
|
||||
|
||||
@ -700,8 +756,8 @@ public:
|
||||
|
||||
if(mem.interpolation == INTERPOLATION_NONE) {
|
||||
/* Data Storage */
|
||||
mem_alloc(mem);
|
||||
mem_copy_to(mem);
|
||||
generic_alloc(mem);
|
||||
generic_copy_to(mem);
|
||||
|
||||
CUdeviceptr cumem;
|
||||
size_t cubytes;
|
||||
@ -891,21 +947,19 @@ public:
|
||||
}
|
||||
else {
|
||||
tex_interp_map.erase(tex_interp_map.find(mem.device_pointer));
|
||||
mem_free(mem);
|
||||
generic_free(mem);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
bool denoising_set_tiles(device_ptr *buffers, DenoisingTask *task)
|
||||
{
|
||||
mem_alloc(task->tiles_mem);
|
||||
|
||||
TilesInfo *tiles = (TilesInfo*) task->tiles_mem.data_pointer;
|
||||
for(int i = 0; i < 9; i++) {
|
||||
tiles->buffers[i] = buffers[i];
|
||||
}
|
||||
|
||||
mem_copy_to(task->tiles_mem);
|
||||
task->tiles_mem.copy_to_device();
|
||||
|
||||
return !have_error();
|
||||
}
|
||||
@ -1272,7 +1326,7 @@ public:
|
||||
task.unmap_neighbor_tiles(rtiles, this);
|
||||
}
|
||||
|
||||
void path_trace(DeviceTask& task, RenderTile& rtile)
|
||||
void path_trace(DeviceTask& task, RenderTile& rtile, device_vector<WorkTile>& work_tiles)
|
||||
{
|
||||
if(have_error())
|
||||
return;
|
||||
@ -1295,8 +1349,7 @@ public:
|
||||
cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
|
||||
|
||||
/* Allocate work tile. */
|
||||
device_vector<WorkTile> work_tiles(this, "work_tiles", MEM_READ_ONLY);
|
||||
work_tiles.resize(1);
|
||||
work_tiles.alloc(1);
|
||||
|
||||
WorkTile *wtile = work_tiles.get_data();
|
||||
wtile->x = rtile.x;
|
||||
@ -1306,9 +1359,6 @@ public:
|
||||
wtile->offset = rtile.offset;
|
||||
wtile->stride = rtile.stride;
|
||||
wtile->buffer = (float*)cuda_device_ptr(rtile.buffer);
|
||||
mem_alloc(work_tiles);
|
||||
|
||||
CUdeviceptr d_work_tiles = cuda_device_ptr(work_tiles.device_pointer);
|
||||
|
||||
/* Prepare work size. More step samples render faster, but for now we
|
||||
* remain conservative for GPUs connected to a display to avoid driver
|
||||
@ -1329,8 +1379,9 @@ public:
|
||||
/* Setup and copy work tile to device. */
|
||||
wtile->start_sample = sample;
|
||||
wtile->num_samples = min(step_samples, end_sample - sample);;
|
||||
mem_copy_to(work_tiles);
|
||||
work_tiles.copy_to_device();
|
||||
|
||||
CUdeviceptr d_work_tiles = cuda_device_ptr(work_tiles.device_pointer);
|
||||
uint total_work_size = wtile->w * wtile->h * wtile->num_samples;
|
||||
uint num_blocks = divide_up(total_work_size, num_threads_per_block);
|
||||
|
||||
@ -1354,8 +1405,6 @@ public:
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
mem_free(work_tiles);
|
||||
}
|
||||
|
||||
void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
|
||||
@ -1508,104 +1557,90 @@ public:
|
||||
|
||||
void pixels_alloc(device_memory& mem)
|
||||
{
|
||||
if(!background) {
|
||||
PixelMem pmem;
|
||||
PixelMem pmem;
|
||||
|
||||
pmem.w = mem.data_width;
|
||||
pmem.h = mem.data_height;
|
||||
pmem.w = mem.data_width;
|
||||
pmem.h = mem.data_height;
|
||||
|
||||
CUDAContextScope scope(this);
|
||||
CUDAContextScope scope(this);
|
||||
|
||||
glGenBuffers(1, &pmem.cuPBO);
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
|
||||
if(mem.data_type == TYPE_HALF)
|
||||
glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(GLhalf)*4, NULL, GL_DYNAMIC_DRAW);
|
||||
else
|
||||
glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(uint8_t)*4, NULL, GL_DYNAMIC_DRAW);
|
||||
glGenBuffers(1, &pmem.cuPBO);
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
|
||||
if(mem.data_type == TYPE_HALF)
|
||||
glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(GLhalf)*4, NULL, GL_DYNAMIC_DRAW);
|
||||
else
|
||||
glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(uint8_t)*4, NULL, GL_DYNAMIC_DRAW);
|
||||
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
|
||||
|
||||
glGenTextures(1, &pmem.cuTexId);
|
||||
glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
|
||||
if(mem.data_type == TYPE_HALF)
|
||||
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, pmem.w, pmem.h, 0, GL_RGBA, GL_HALF_FLOAT, NULL);
|
||||
else
|
||||
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, pmem.w, pmem.h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
|
||||
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
|
||||
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
|
||||
glBindTexture(GL_TEXTURE_2D, 0);
|
||||
glGenTextures(1, &pmem.cuTexId);
|
||||
glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
|
||||
if(mem.data_type == TYPE_HALF)
|
||||
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, pmem.w, pmem.h, 0, GL_RGBA, GL_HALF_FLOAT, NULL);
|
||||
else
|
||||
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, pmem.w, pmem.h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
|
||||
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
|
||||
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
|
||||
glBindTexture(GL_TEXTURE_2D, 0);
|
||||
|
||||
CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE);
|
||||
CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE);
|
||||
|
||||
if(result == CUDA_SUCCESS) {
|
||||
mem.device_pointer = pmem.cuTexId;
|
||||
pixel_mem_map[mem.device_pointer] = pmem;
|
||||
if(result == CUDA_SUCCESS) {
|
||||
mem.device_pointer = pmem.cuTexId;
|
||||
pixel_mem_map[mem.device_pointer] = pmem;
|
||||
|
||||
mem.device_size = mem.memory_size();
|
||||
stats.mem_alloc(mem.device_size);
|
||||
mem.device_size = mem.memory_size();
|
||||
stats.mem_alloc(mem.device_size);
|
||||
|
||||
return;
|
||||
}
|
||||
else {
|
||||
/* failed to register buffer, fallback to no interop */
|
||||
glDeleteBuffers(1, &pmem.cuPBO);
|
||||
glDeleteTextures(1, &pmem.cuTexId);
|
||||
|
||||
background = true;
|
||||
}
|
||||
return;
|
||||
}
|
||||
else {
|
||||
/* failed to register buffer, fallback to no interop */
|
||||
glDeleteBuffers(1, &pmem.cuPBO);
|
||||
glDeleteTextures(1, &pmem.cuTexId);
|
||||
|
||||
Device::pixels_alloc(mem);
|
||||
background = true;
|
||||
}
|
||||
}
|
||||
|
||||
void pixels_copy_from(device_memory& mem, int y, int w, int h)
|
||||
{
|
||||
if(!background) {
|
||||
PixelMem pmem = pixel_mem_map[mem.device_pointer];
|
||||
PixelMem pmem = pixel_mem_map[mem.device_pointer];
|
||||
|
||||
CUDAContextScope scope(this);
|
||||
CUDAContextScope scope(this);
|
||||
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
|
||||
uchar *pixels = (uchar*)glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_ONLY);
|
||||
size_t offset = sizeof(uchar)*4*y*w;
|
||||
memcpy((uchar*)mem.data_pointer + offset, pixels + offset, sizeof(uchar)*4*w*h);
|
||||
glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
|
||||
|
||||
return;
|
||||
}
|
||||
|
||||
Device::pixels_copy_from(mem, y, w, h);
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
|
||||
uchar *pixels = (uchar*)glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_ONLY);
|
||||
size_t offset = sizeof(uchar)*4*y*w;
|
||||
memcpy((uchar*)mem.data_pointer + offset, pixels + offset, sizeof(uchar)*4*w*h);
|
||||
glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
|
||||
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
|
||||
}
|
||||
|
||||
void pixels_free(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
if(!background) {
|
||||
PixelMem pmem = pixel_mem_map[mem.device_pointer];
|
||||
PixelMem pmem = pixel_mem_map[mem.device_pointer];
|
||||
|
||||
CUDAContextScope scope(this);
|
||||
CUDAContextScope scope(this);
|
||||
|
||||
cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource));
|
||||
glDeleteBuffers(1, &pmem.cuPBO);
|
||||
glDeleteTextures(1, &pmem.cuTexId);
|
||||
cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource));
|
||||
glDeleteBuffers(1, &pmem.cuPBO);
|
||||
glDeleteTextures(1, &pmem.cuTexId);
|
||||
|
||||
pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer));
|
||||
mem.device_pointer = 0;
|
||||
pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer));
|
||||
mem.device_pointer = 0;
|
||||
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_size = 0;
|
||||
|
||||
return;
|
||||
}
|
||||
|
||||
Device::pixels_free(mem);
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_size = 0;
|
||||
}
|
||||
}
|
||||
|
||||
void draw_pixels(device_memory& mem, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
|
||||
const DeviceDrawParams &draw_params)
|
||||
{
|
||||
assert(mem.type == MEM_PIXELS);
|
||||
|
||||
if(!background) {
|
||||
PixelMem pmem = pixel_mem_map[mem.device_pointer];
|
||||
float *vpointer;
|
||||
@ -1724,6 +1759,8 @@ public:
|
||||
}
|
||||
}
|
||||
|
||||
device_vector<WorkTile> work_tiles(this, "work_tiles", MEM_READ_ONLY);
|
||||
|
||||
/* keep rendering tiles until done */
|
||||
while(task->acquire_tile(this, tile)) {
|
||||
if(tile.task == RenderTile::PATH_TRACE) {
|
||||
@ -1732,7 +1769,7 @@ public:
|
||||
split_kernel->path_trace(task, tile, void_buffer, void_buffer);
|
||||
}
|
||||
else {
|
||||
path_trace(*task, tile);
|
||||
path_trace(*task, tile, work_tiles);
|
||||
}
|
||||
}
|
||||
else if(tile.task == RenderTile::DENOISE) {
|
||||
@ -1750,6 +1787,8 @@ public:
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
work_tiles.free();
|
||||
}
|
||||
else if(task->type == DeviceTask::SHADER) {
|
||||
shader(*task);
|
||||
@ -1884,8 +1923,8 @@ uint64_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory
|
||||
CUDAContextScope scope(device);
|
||||
|
||||
device_vector<uint64_t> size_buffer(device, "size_buffer", MEM_READ_WRITE);
|
||||
size_buffer.resize(1);
|
||||
device->mem_alloc(size_buffer);
|
||||
size_buffer.alloc(1);
|
||||
size_buffer.zero_to_device();
|
||||
|
||||
uint threads = num_threads;
|
||||
CUdeviceptr d_size = device->cuda_device_ptr(size_buffer.device_pointer);
|
||||
@ -1908,9 +1947,9 @@ uint64_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory
|
||||
1, 1, 1,
|
||||
0, 0, (void**)&args, 0));
|
||||
|
||||
device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t));
|
||||
size_buffer.copy_from_device(0, 1, 1);
|
||||
size_t size = size_buffer[0];
|
||||
device->mem_free(size_buffer);
|
||||
size_buffer.free();
|
||||
|
||||
return size;
|
||||
}
|
||||
|
@ -44,7 +44,7 @@ void DenoisingTask::init_from_devicetask(const DeviceTask &task)
|
||||
|
||||
void DenoisingTask::tiles_from_rendertiles(RenderTile *rtiles)
|
||||
{
|
||||
tiles = (TilesInfo*) tiles_mem.resize(sizeof(TilesInfo)/sizeof(int));
|
||||
tiles = (TilesInfo*) tiles_mem.alloc(sizeof(TilesInfo)/sizeof(int));
|
||||
|
||||
device_ptr buffers[9];
|
||||
for(int i = 0; i < 9; i++) {
|
||||
@ -75,8 +75,7 @@ bool DenoisingTask::run_denoising()
|
||||
buffer.w = align_up(rect.z - rect.x, 4);
|
||||
buffer.h = rect.w - rect.y;
|
||||
buffer.pass_stride = align_up(buffer.w * buffer.h, divide_up(device->mem_address_alignment(), sizeof(float)));
|
||||
buffer.mem.resize(buffer.pass_stride * buffer.passes);
|
||||
device->mem_alloc(buffer.mem);
|
||||
buffer.mem.alloc_to_device(buffer.pass_stride * buffer.passes);
|
||||
|
||||
device_ptr null_ptr = (device_ptr) 0;
|
||||
|
||||
@ -161,8 +160,7 @@ bool DenoisingTask::run_denoising()
|
||||
int num_color_passes = 3;
|
||||
|
||||
device_only_memory<float> temp_color(device, "Denoising temporary color");
|
||||
temp_color.resize(3*buffer.pass_stride);
|
||||
device->mem_alloc(temp_color);
|
||||
temp_color.alloc_to_device(3*buffer.pass_stride);
|
||||
|
||||
for(int pass = 0; pass < num_color_passes; pass++) {
|
||||
device_sub_ptr color_pass(temp_color, pass*buffer.pass_stride, buffer.pass_stride);
|
||||
@ -177,31 +175,25 @@ bool DenoisingTask::run_denoising()
|
||||
functions.detect_outliers(temp_color.device_pointer, *color_var_pass, *depth_pass, *output_pass);
|
||||
}
|
||||
|
||||
device->mem_free(temp_color);
|
||||
temp_color.free();
|
||||
}
|
||||
|
||||
storage.w = filter_area.z;
|
||||
storage.h = filter_area.w;
|
||||
storage.transform.resize(storage.w*storage.h*TRANSFORM_SIZE);
|
||||
storage.rank.resize(storage.w*storage.h);
|
||||
device->mem_alloc(storage.transform);
|
||||
device->mem_alloc(storage.rank);
|
||||
storage.transform.alloc_to_device(storage.w*storage.h*TRANSFORM_SIZE);
|
||||
storage.rank.alloc_to_device(storage.w*storage.h);
|
||||
|
||||
functions.construct_transform();
|
||||
|
||||
device_only_memory<float> temporary_1(device, "Denoising NLM temporary 1");
|
||||
device_only_memory<float> temporary_2(device, "Denoising NLM temporary 2");
|
||||
temporary_1.resize(buffer.w*buffer.h);
|
||||
temporary_2.resize(buffer.w*buffer.h);
|
||||
device->mem_alloc(temporary_1);
|
||||
device->mem_alloc(temporary_2);
|
||||
temporary_1.alloc_to_device(buffer.w*buffer.h);
|
||||
temporary_2.alloc_to_device(buffer.w*buffer.h);
|
||||
reconstruction_state.temporary_1_ptr = temporary_1.device_pointer;
|
||||
reconstruction_state.temporary_2_ptr = temporary_2.device_pointer;
|
||||
|
||||
storage.XtWX.resize(storage.w*storage.h*XTWX_SIZE);
|
||||
storage.XtWY.resize(storage.w*storage.h*XTWY_SIZE);
|
||||
device->mem_alloc(storage.XtWX);
|
||||
device->mem_alloc(storage.XtWY);
|
||||
storage.XtWX.alloc_to_device(storage.w*storage.h*XTWX_SIZE);
|
||||
storage.XtWY.alloc_to_device(storage.w*storage.h*XTWY_SIZE);
|
||||
|
||||
reconstruction_state.filter_rect = make_int4(filter_area.x-rect.x, filter_area.y-rect.y, storage.w, storage.h);
|
||||
int tile_coordinate_offset = filter_area.y*render_buffer.stride + filter_area.x;
|
||||
@ -218,14 +210,14 @@ bool DenoisingTask::run_denoising()
|
||||
functions.reconstruct(*color_ptr, *color_var_ptr, render_buffer.ptr);
|
||||
}
|
||||
|
||||
device->mem_free(storage.XtWX);
|
||||
device->mem_free(storage.XtWY);
|
||||
device->mem_free(storage.transform);
|
||||
device->mem_free(storage.rank);
|
||||
device->mem_free(temporary_1);
|
||||
device->mem_free(temporary_2);
|
||||
device->mem_free(buffer.mem);
|
||||
device->mem_free(tiles_mem);
|
||||
storage.XtWX.free();
|
||||
storage.XtWY.free();
|
||||
storage.transform.free();
|
||||
storage.rank.free();
|
||||
temporary_1.free();
|
||||
temporary_2.free();
|
||||
buffer.mem.free();
|
||||
tiles_mem.free();
|
||||
return true;
|
||||
}
|
||||
|
||||
|
@ -43,6 +43,68 @@ device_memory::~device_memory()
|
||||
{
|
||||
}
|
||||
|
||||
device_ptr device_memory::host_alloc(size_t size)
|
||||
{
|
||||
if(!size) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
size_t alignment = device->mem_address_alignment();
|
||||
device_ptr ptr = (device_ptr)util_aligned_malloc(size, alignment);
|
||||
|
||||
if(ptr) {
|
||||
util_guarded_mem_alloc(size);
|
||||
}
|
||||
else {
|
||||
throw std::bad_alloc();
|
||||
}
|
||||
|
||||
return ptr;
|
||||
}
|
||||
|
||||
void device_memory::host_free(device_ptr ptr, size_t size)
|
||||
{
|
||||
if(ptr) {
|
||||
util_guarded_mem_free(size);
|
||||
util_aligned_free((void*)ptr);
|
||||
}
|
||||
}
|
||||
|
||||
void device_memory::device_alloc()
|
||||
{
|
||||
assert(!device_pointer && type != MEM_TEXTURE);
|
||||
device->mem_alloc(*this);
|
||||
}
|
||||
|
||||
void device_memory::device_free()
|
||||
{
|
||||
if(device_pointer) {
|
||||
device->mem_free(*this);
|
||||
}
|
||||
}
|
||||
|
||||
void device_memory::device_copy_to()
|
||||
{
|
||||
assert(type != MEM_PIXELS && type != MEM_WRITE_ONLY);
|
||||
if(data_size) {
|
||||
device->mem_copy_to(*this);
|
||||
}
|
||||
}
|
||||
|
||||
void device_memory::device_copy_from(int y, int w, int h, int elem)
|
||||
{
|
||||
assert(type != MEM_TEXTURE && type != MEM_READ_ONLY);
|
||||
device->mem_copy_from(*this, y, w, h, elem);
|
||||
}
|
||||
|
||||
void device_memory::device_zero()
|
||||
{
|
||||
assert(type != MEM_PIXELS && type != MEM_WRITE_ONLY);
|
||||
if(data_size) {
|
||||
device->mem_zero(*this);
|
||||
}
|
||||
}
|
||||
|
||||
/* Device Sub Ptr */
|
||||
|
||||
device_sub_ptr::device_sub_ptr(device_memory& mem, int offset, int size)
|
||||
|
@ -19,14 +19,7 @@
|
||||
|
||||
/* Device Memory
|
||||
*
|
||||
* This file defines data types that can be used in device memory arrays, and
|
||||
* a device_vector<T> type to store such arrays.
|
||||
*
|
||||
* device_vector<T> contains an STL vector, metadata about the data type,
|
||||
* dimensions, elements, and a device pointer. For the CPU device this is just
|
||||
* a pointer to the STL vector data, as no copying needs to take place. For
|
||||
* other devices this is a pointer to device memory, where we will copy memory
|
||||
* to and from. */
|
||||
* Data types for allocating, copying and freeing device memory. */
|
||||
|
||||
#include "util/util_debug.h"
|
||||
#include "util/util_half.h"
|
||||
@ -41,7 +34,9 @@ class Device;
|
||||
enum MemoryType {
|
||||
MEM_READ_ONLY,
|
||||
MEM_WRITE_ONLY,
|
||||
MEM_READ_WRITE
|
||||
MEM_READ_WRITE,
|
||||
MEM_TEXTURE,
|
||||
MEM_PIXELS
|
||||
};
|
||||
|
||||
/* Supported Data Types */
|
||||
@ -172,7 +167,10 @@ template<> struct device_type_traits<uint64_t> {
|
||||
static const int num_elements = 1;
|
||||
};
|
||||
|
||||
/* Device Memory */
|
||||
/* Device Memory
|
||||
*
|
||||
* Base class for all device memory. This should not be allocated directly,
|
||||
* instead the appropriate subclass can be used. */
|
||||
|
||||
class device_memory
|
||||
{
|
||||
@ -182,7 +180,7 @@ public:
|
||||
return elements*data_elements*datatype_size(data_type);
|
||||
}
|
||||
|
||||
/* data information */
|
||||
/* Data information. */
|
||||
DataType data_type;
|
||||
int data_elements;
|
||||
device_ptr data_pointer;
|
||||
@ -196,25 +194,39 @@ public:
|
||||
InterpolationType interpolation;
|
||||
ExtensionType extension;
|
||||
|
||||
/* device pointer */
|
||||
/* Device pointer. */
|
||||
Device *device;
|
||||
device_ptr device_pointer;
|
||||
|
||||
device_memory(Device *device, const char *name, MemoryType type);
|
||||
virtual ~device_memory();
|
||||
|
||||
void resize(size_t size)
|
||||
{
|
||||
data_size = size;
|
||||
data_width = size;
|
||||
}
|
||||
|
||||
protected:
|
||||
/* no copying */
|
||||
/* Only create through subclasses. */
|
||||
device_memory(Device *device, const char *name, MemoryType type);
|
||||
|
||||
/* No copying allowed. */
|
||||
device_memory(const device_memory&);
|
||||
device_memory& operator = (const device_memory&);
|
||||
|
||||
/* Host allocation on the device. All data_pointer memory should be
|
||||
* allocated with these functions, for devices that support using
|
||||
* the same pointer for host and device. */
|
||||
device_ptr host_alloc(size_t size);
|
||||
void host_free(device_ptr ptr, size_t size);
|
||||
|
||||
/* Device memory allocation and copying. */
|
||||
void device_alloc();
|
||||
void device_free();
|
||||
void device_copy_to();
|
||||
void device_copy_from(int y, int w, int h, int elem);
|
||||
void device_zero();
|
||||
};
|
||||
|
||||
/* Device Only Memory
|
||||
*
|
||||
* Working memory only needed by the device, with no corresponding allocation
|
||||
* on the host. Only used internally in the device implementations. */
|
||||
|
||||
template<typename T>
|
||||
class device_only_memory : public device_memory
|
||||
{
|
||||
@ -226,18 +238,43 @@ public:
|
||||
data_elements = max(device_type_traits<T>::num_elements, 1);
|
||||
}
|
||||
|
||||
void resize(size_t num)
|
||||
virtual ~device_only_memory()
|
||||
{
|
||||
device_memory::resize(num*sizeof(T));
|
||||
free();
|
||||
}
|
||||
|
||||
void alloc_to_device(size_t num)
|
||||
{
|
||||
data_size = num*sizeof(T);
|
||||
device_alloc();
|
||||
}
|
||||
|
||||
void free()
|
||||
{
|
||||
device_free();
|
||||
}
|
||||
|
||||
void zero_to_device()
|
||||
{
|
||||
device_zero();
|
||||
}
|
||||
};
|
||||
|
||||
/* Device Vector */
|
||||
/* Device Vector
|
||||
*
|
||||
* Data vector to exchange data between host and device. Memory will be
|
||||
* allocated on the host first with alloc() and resize, and then filled
|
||||
* in and copied to the device with copy_to_device(). Or alternatively
|
||||
* allocated and set to zero on the device with zero_to_device().
|
||||
*
|
||||
* When using memory type MEM_TEXTURE, a pointer to this memory will be
|
||||
* automatically attached to kernel globals, using the provided name
|
||||
* matching an entry in kernel_textures.h. */
|
||||
|
||||
template<typename T> class device_vector : public device_memory
|
||||
{
|
||||
public:
|
||||
device_vector(Device *device, const char *name, MemoryType type = MEM_READ_ONLY)
|
||||
device_vector(Device *device, const char *name, MemoryType type)
|
||||
: device_memory(device, name, type)
|
||||
{
|
||||
data_type = device_type_traits<T>::data_type;
|
||||
@ -246,84 +283,175 @@ public:
|
||||
assert(data_elements > 0);
|
||||
}
|
||||
|
||||
virtual ~device_vector() {}
|
||||
|
||||
/* vector functions */
|
||||
T *resize(size_t width, size_t height = 0, size_t depth = 0)
|
||||
virtual ~device_vector()
|
||||
{
|
||||
data_size = width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth);
|
||||
if(data.resize(data_size) == NULL) {
|
||||
clear();
|
||||
return NULL;
|
||||
free();
|
||||
}
|
||||
|
||||
/* Host memory allocation. */
|
||||
T *alloc(size_t width, size_t height = 0, size_t depth = 0)
|
||||
{
|
||||
size_t new_size = size(width, height, depth);
|
||||
|
||||
if(new_size != data_size) {
|
||||
device_free();
|
||||
host_free(data_pointer, sizeof(T)*data_size);
|
||||
data_pointer = host_alloc(sizeof(T)*new_size);
|
||||
}
|
||||
|
||||
data_size = new_size;
|
||||
data_width = width;
|
||||
data_height = height;
|
||||
data_depth = depth;
|
||||
if(data_size == 0) {
|
||||
data_pointer = 0;
|
||||
return NULL;
|
||||
}
|
||||
data_pointer = (device_ptr)&data[0];
|
||||
return &data[0];
|
||||
assert(device_ptr == 0);
|
||||
|
||||
return get_data();
|
||||
}
|
||||
|
||||
/* Host memory resize. Only use this if the original data needs to be
|
||||
* preserved, it is faster to call alloc() if it can be discarded. */
|
||||
T *resize(size_t width, size_t height = 0, size_t depth = 0)
|
||||
{
|
||||
size_t new_size = size(width, height, depth);
|
||||
|
||||
if(new_size != data_size) {
|
||||
device_ptr new_ptr = host_alloc(sizeof(T)*new_size);
|
||||
|
||||
if(new_size && data_size) {
|
||||
size_t min_size = ((new_size < data_size)? new_size: data_size);
|
||||
memcpy((T*)new_ptr, (T*)data_pointer, sizeof(T)*min_size);
|
||||
}
|
||||
|
||||
device_free();
|
||||
host_free(data_pointer, sizeof(T)*data_size);
|
||||
data_pointer = new_ptr;
|
||||
}
|
||||
|
||||
data_size = new_size;
|
||||
data_width = width;
|
||||
data_height = height;
|
||||
data_depth = depth;
|
||||
assert(device_ptr == 0);
|
||||
|
||||
return get_data();
|
||||
}
|
||||
|
||||
/* Take over data from an existing array. */
|
||||
void steal_data(array<T>& from)
|
||||
{
|
||||
data.steal_data(from);
|
||||
data_size = data.size();
|
||||
data_pointer = (data_size)? (device_ptr)&data[0]: 0;
|
||||
data_width = data_size;
|
||||
data_height = 0;
|
||||
data_depth = 0;
|
||||
}
|
||||
device_free();
|
||||
host_free(data_pointer, sizeof(T)*data_size);
|
||||
|
||||
void clear()
|
||||
{
|
||||
data.clear();
|
||||
data_pointer = 0;
|
||||
data_size = from.size();
|
||||
data_width = 0;
|
||||
data_height = 0;
|
||||
data_depth = 0;
|
||||
data_pointer = (device_ptr)from.steal_pointer();
|
||||
assert(device_pointer == 0);
|
||||
}
|
||||
|
||||
/* Free device and host memory. */
|
||||
void free()
|
||||
{
|
||||
device_free();
|
||||
host_free(data_pointer, sizeof(T)*data_size);
|
||||
|
||||
data_size = 0;
|
||||
device_pointer = 0;
|
||||
data_width = 0;
|
||||
data_height = 0;
|
||||
data_depth = 0;
|
||||
data_pointer = 0;
|
||||
assert(device_pointer == 0);
|
||||
}
|
||||
|
||||
size_t size()
|
||||
{
|
||||
return data.size();
|
||||
return data_size;
|
||||
}
|
||||
|
||||
T* get_data()
|
||||
{
|
||||
return &data[0];
|
||||
return (T*)data_pointer;
|
||||
}
|
||||
|
||||
T& operator[](size_t i)
|
||||
{
|
||||
return data[i];
|
||||
assert(i < data_size);
|
||||
return get_data()[i];
|
||||
}
|
||||
|
||||
private:
|
||||
array<T> data;
|
||||
void copy_to_device()
|
||||
{
|
||||
device_copy_to();
|
||||
}
|
||||
|
||||
void copy_from_device(int y, int w, int h)
|
||||
{
|
||||
device_copy_from(y, w, h, sizeof(T));
|
||||
}
|
||||
|
||||
void zero_to_device()
|
||||
{
|
||||
device_zero();
|
||||
}
|
||||
|
||||
protected:
|
||||
size_t size(size_t width, size_t height, size_t depth)
|
||||
{
|
||||
return width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth);
|
||||
}
|
||||
};
|
||||
|
||||
/* A device_sub_ptr is a pointer into another existing memory.
|
||||
* Therefore, it is not allocated separately, but just created from the already allocated base memory.
|
||||
* It is freed automatically when it goes out of scope, which should happen before the base memory is freed.
|
||||
* Note that some devices require the offset and size of the sub_ptr to be properly aligned. */
|
||||
/* Pixel Memory
|
||||
*
|
||||
* Device memory to efficiently draw as pixels to the screen in interactive
|
||||
* rendering. Only copying pixels from the device is supported, not copying to. */
|
||||
|
||||
template<typename T> class device_pixels : public device_vector<T>
|
||||
{
|
||||
public:
|
||||
device_pixels(Device *device, const char *name)
|
||||
: device_vector<T>(device, name, MEM_PIXELS)
|
||||
{
|
||||
}
|
||||
|
||||
void alloc_to_device(size_t width, size_t height, size_t depth = 0)
|
||||
{
|
||||
device_vector<T>::alloc(width, height, depth);
|
||||
device_memory::device_alloc();
|
||||
}
|
||||
|
||||
T *copy_from_device(int y, int w, int h)
|
||||
{
|
||||
device_memory::device_copy_from(y, w, h, sizeof(T));
|
||||
return device_vector<T>::get_data();
|
||||
}
|
||||
};
|
||||
|
||||
/* Device Sub Memory
|
||||
*
|
||||
* Pointer into existing memory. It is not allocated separately, but created
|
||||
* from an already allocated base memory. It is freed automatically when it
|
||||
* goes out of scope, which should happen before base memory is freed.
|
||||
*
|
||||
* Note: some devices require offset and size of the sub_ptr to be properly
|
||||
* aligned to device->mem_address_alingment(). */
|
||||
|
||||
class device_sub_ptr
|
||||
{
|
||||
public:
|
||||
device_sub_ptr(device_memory& mem, int offset, int size);
|
||||
~device_sub_ptr();
|
||||
/* No copying. */
|
||||
device_sub_ptr& operator = (const device_sub_ptr&);
|
||||
|
||||
device_ptr operator*() const
|
||||
{
|
||||
return ptr;
|
||||
}
|
||||
|
||||
protected:
|
||||
/* No copying. */
|
||||
device_sub_ptr& operator = (const device_sub_ptr&);
|
||||
|
||||
Device *device;
|
||||
device_ptr ptr;
|
||||
};
|
||||
|
@ -43,10 +43,10 @@ public:
|
||||
};
|
||||
|
||||
list<SubDevice> devices;
|
||||
device_ptr unique_ptr;
|
||||
device_ptr unique_key;
|
||||
|
||||
MultiDevice(DeviceInfo& info, Stats &stats, bool background_)
|
||||
: Device(info, stats, background_), unique_ptr(1)
|
||||
: Device(info, stats, background_), unique_key(1)
|
||||
{
|
||||
Device *device;
|
||||
|
||||
@ -108,68 +108,87 @@ public:
|
||||
|
||||
void mem_alloc(device_memory& mem)
|
||||
{
|
||||
device_ptr key = unique_key++;
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device = sub.device;
|
||||
mem.device_pointer = 0;
|
||||
|
||||
sub.device->mem_alloc(mem);
|
||||
sub.ptr_map[unique_ptr] = mem.device_pointer;
|
||||
sub.ptr_map[key] = mem.device_pointer;
|
||||
}
|
||||
|
||||
mem.device_pointer = unique_ptr++;
|
||||
stats.mem_alloc(mem.device_size);
|
||||
mem.device = this;
|
||||
mem.device_pointer = key;
|
||||
}
|
||||
|
||||
void mem_copy_to(device_memory& mem)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
device_ptr existing_key = mem.device_pointer;
|
||||
device_ptr key = (existing_key)? existing_key: unique_key++;
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
mem.device = sub.device;
|
||||
mem.device_pointer = (existing_key)? sub.ptr_map[existing_key]: 0;
|
||||
|
||||
sub.device->mem_copy_to(mem);
|
||||
sub.ptr_map[key] = mem.device_pointer;
|
||||
}
|
||||
|
||||
mem.device_pointer = tmp;
|
||||
mem.device = this;
|
||||
mem.device_pointer = key;
|
||||
}
|
||||
|
||||
void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
device_ptr key = mem.device_pointer;
|
||||
int i = 0, sub_h = h/devices.size();
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
int sy = y + i*sub_h;
|
||||
int sh = (i == (int)devices.size() - 1)? h - sub_h*i: sub_h;
|
||||
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
mem.device = sub.device;
|
||||
mem.device_pointer = sub.ptr_map[key];
|
||||
|
||||
sub.device->mem_copy_from(mem, sy, w, sh, elem);
|
||||
i++;
|
||||
}
|
||||
|
||||
mem.device_pointer = tmp;
|
||||
mem.device = this;
|
||||
mem.device_pointer = key;
|
||||
}
|
||||
|
||||
void mem_zero(device_memory& mem)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
device_ptr existing_key = mem.device_pointer;
|
||||
device_ptr key = (existing_key)? existing_key: unique_key++;
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
mem.device = sub.device;
|
||||
mem.device_pointer = (existing_key)? sub.ptr_map[existing_key]: 0;
|
||||
|
||||
sub.device->mem_zero(mem);
|
||||
sub.ptr_map[key] = mem.device_pointer;
|
||||
}
|
||||
|
||||
mem.device_pointer = tmp;
|
||||
mem.device = this;
|
||||
mem.device_pointer = key;
|
||||
}
|
||||
|
||||
void mem_free(device_memory& mem)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
stats.mem_free(mem.device_size);
|
||||
device_ptr key = mem.device_pointer;
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
mem.device = sub.device;
|
||||
mem.device_pointer = sub.ptr_map[key];
|
||||
|
||||
sub.device->mem_free(mem);
|
||||
sub.ptr_map.erase(sub.ptr_map.find(tmp));
|
||||
sub.ptr_map.erase(sub.ptr_map.find(key));
|
||||
}
|
||||
|
||||
mem.device = this;
|
||||
mem.device_pointer = 0;
|
||||
}
|
||||
|
||||
@ -179,81 +198,10 @@ public:
|
||||
sub.device->const_copy_to(name, host, size);
|
||||
}
|
||||
|
||||
void tex_alloc(device_memory& mem)
|
||||
{
|
||||
VLOG(1) << "Texture allocate: " << mem.name << ", "
|
||||
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
|
||||
<< string_human_readable_size(mem.memory_size()) << ")";
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = 0;
|
||||
sub.device->tex_alloc(mem);
|
||||
sub.ptr_map[unique_ptr] = mem.device_pointer;
|
||||
}
|
||||
|
||||
mem.device_pointer = unique_ptr++;
|
||||
stats.mem_alloc(mem.device_size);
|
||||
}
|
||||
|
||||
void tex_free(device_memory& mem)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
stats.mem_free(mem.device_size);
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
sub.device->tex_free(mem);
|
||||
sub.ptr_map.erase(sub.ptr_map.find(tmp));
|
||||
}
|
||||
|
||||
mem.device_pointer = 0;
|
||||
}
|
||||
|
||||
void pixels_alloc(device_memory& mem)
|
||||
{
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = 0;
|
||||
sub.device->pixels_alloc(mem);
|
||||
sub.ptr_map[unique_ptr] = mem.device_pointer;
|
||||
}
|
||||
|
||||
mem.device_pointer = unique_ptr++;
|
||||
}
|
||||
|
||||
void pixels_free(device_memory& mem)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
sub.device->pixels_free(mem);
|
||||
sub.ptr_map.erase(sub.ptr_map.find(tmp));
|
||||
}
|
||||
|
||||
mem.device_pointer = 0;
|
||||
}
|
||||
|
||||
void pixels_copy_from(device_memory& mem, int y, int w, int h)
|
||||
{
|
||||
device_ptr tmp = mem.device_pointer;
|
||||
int i = 0, sub_h = h/devices.size();
|
||||
|
||||
foreach(SubDevice& sub, devices) {
|
||||
int sy = y + i*sub_h;
|
||||
int sh = (i == (int)devices.size() - 1)? h - sub_h*i: sub_h;
|
||||
|
||||
mem.device_pointer = sub.ptr_map[tmp];
|
||||
sub.device->pixels_copy_from(mem, sy, w, sh);
|
||||
i++;
|
||||
}
|
||||
|
||||
mem.device_pointer = tmp;
|
||||
}
|
||||
|
||||
void draw_pixels(device_memory& rgba, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
|
||||
const DeviceDrawParams &draw_params)
|
||||
{
|
||||
device_ptr tmp = rgba.device_pointer;
|
||||
device_ptr key = rgba.device_pointer;
|
||||
int i = 0, sub_h = h/devices.size();
|
||||
int sub_height = height/devices.size();
|
||||
|
||||
@ -264,12 +212,12 @@ public:
|
||||
int sdy = dy + i*sub_height;
|
||||
/* adjust math for w/width */
|
||||
|
||||
rgba.device_pointer = sub.ptr_map[tmp];
|
||||
rgba.device_pointer = sub.ptr_map[key];
|
||||
sub.device->draw_pixels(rgba, sy, w, sh, dx, sdy, width, sheight, transparent, draw_params);
|
||||
i++;
|
||||
}
|
||||
|
||||
rgba.device_pointer = tmp;
|
||||
rgba.device_pointer = key;
|
||||
}
|
||||
|
||||
void map_tile(Device *sub_device, RenderTile& tile)
|
||||
@ -304,15 +252,21 @@ public:
|
||||
* to the current device now, for the duration of the denoising task.
|
||||
* Note that this temporarily modifies the RenderBuffers and calls
|
||||
* the device, so this function is not thread safe. */
|
||||
if(tiles[i].buffers->device != sub_device) {
|
||||
device_vector<float> &mem = tiles[i].buffers->buffer;
|
||||
|
||||
device_vector<float> &mem = tiles[i].buffers->buffer;
|
||||
if(mem.device != sub_device) {
|
||||
tiles[i].buffers->copy_from_device();
|
||||
|
||||
Device *original_device = mem.device;
|
||||
device_ptr original_ptr = mem.device_pointer;
|
||||
|
||||
mem.device = sub_device;
|
||||
mem.device_pointer = 0;
|
||||
|
||||
sub_device->mem_alloc(mem);
|
||||
sub_device->mem_copy_to(mem);
|
||||
tiles[i].buffer = mem.device_pointer;
|
||||
|
||||
mem.device = original_device;
|
||||
mem.device_pointer = original_ptr;
|
||||
}
|
||||
}
|
||||
@ -324,25 +278,30 @@ public:
|
||||
if(!tiles[i].buffers) {
|
||||
continue;
|
||||
}
|
||||
if(tiles[i].buffers->device != sub_device) {
|
||||
device_vector<float> &mem = tiles[i].buffers->buffer;
|
||||
|
||||
device_vector<float> &mem = tiles[i].buffers->buffer;
|
||||
if(mem.device != sub_device) {
|
||||
Device *original_device = mem.device;
|
||||
device_ptr original_ptr = mem.device_pointer;
|
||||
size_t original_size = mem.device_size;
|
||||
|
||||
mem.device = sub_device;
|
||||
mem.device_pointer = tiles[i].buffer;
|
||||
|
||||
/* Copy denoised tile to the host. */
|
||||
if(i == 4) {
|
||||
tiles[i].buffers->copy_from_device(sub_device);
|
||||
tiles[i].buffers->copy_from_device();
|
||||
}
|
||||
|
||||
size_t mem_size = mem.device_size;
|
||||
sub_device->mem_free(mem);
|
||||
|
||||
mem.device = original_device;
|
||||
mem.device_pointer = original_ptr;
|
||||
mem.device_size = mem_size;
|
||||
mem.device_size = original_size;
|
||||
|
||||
/* Copy denoised tile to the original device. */
|
||||
if(i == 4) {
|
||||
tiles[i].buffers->device->mem_copy_to(mem);
|
||||
mem.copy_to_device();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
@ -172,36 +172,6 @@ public:
|
||||
snd.write_buffer(host, size);
|
||||
}
|
||||
|
||||
void tex_alloc(device_memory& mem)
|
||||
{
|
||||
VLOG(1) << "Texture allocate: " << mem.name << ", "
|
||||
<< string_human_readable_number(mem.memory_size()) << " bytes. ("
|
||||
<< string_human_readable_size(mem.memory_size()) << ")";
|
||||
|
||||
thread_scoped_lock lock(rpc_lock);
|
||||
|
||||
mem.device_pointer = ++mem_counter;
|
||||
|
||||
RPCSend snd(socket, &error_func, "tex_alloc");
|
||||
snd.add(mem);
|
||||
snd.write();
|
||||
snd.write_buffer((void*)mem.data_pointer, mem.memory_size());
|
||||
}
|
||||
|
||||
void tex_free(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
thread_scoped_lock lock(rpc_lock);
|
||||
|
||||
RPCSend snd(socket, &error_func, "tex_free");
|
||||
|
||||
snd.add(mem);
|
||||
snd.write();
|
||||
|
||||
mem.device_pointer = 0;
|
||||
}
|
||||
}
|
||||
|
||||
bool load_kernels(const DeviceRequestedFeatures& requested_features)
|
||||
{
|
||||
if(error_func.have_error())
|
||||
@ -310,7 +280,7 @@ public:
|
||||
snd.write();
|
||||
}
|
||||
|
||||
int get_split_task_count(DeviceTask& task)
|
||||
int get_split_task_count(DeviceTask&)
|
||||
{
|
||||
return 1;
|
||||
}
|
||||
@ -464,21 +434,17 @@ protected:
|
||||
rcv.read(mem, name);
|
||||
lock.unlock();
|
||||
|
||||
/* Allocate host side data buffer. */
|
||||
size_t data_size = mem.memory_size();
|
||||
device_ptr client_pointer = mem.device_pointer;
|
||||
|
||||
/* create a memory buffer for the device buffer */
|
||||
size_t data_size = mem.memory_size();
|
||||
DataVector &data_v = data_vector_insert(client_pointer, data_size);
|
||||
mem.data_pointer = (data_size)? (device_ptr)&(data_v[0]): 0;
|
||||
|
||||
if(data_size)
|
||||
mem.data_pointer = (device_ptr)&(data_v[0]);
|
||||
else
|
||||
mem.data_pointer = 0;
|
||||
|
||||
/* perform the allocation on the actual device */
|
||||
/* Perform the allocation on the actual device. */
|
||||
device->mem_alloc(mem);
|
||||
|
||||
/* store a mapping to/from client_pointer and real device pointer */
|
||||
/* Store a mapping to/from client_pointer and real device pointer. */
|
||||
pointer_mapping_insert(client_pointer, mem.device_pointer);
|
||||
}
|
||||
else if(rcv.name == "mem_copy_to") {
|
||||
@ -487,23 +453,33 @@ protected:
|
||||
rcv.read(mem, name);
|
||||
lock.unlock();
|
||||
|
||||
size_t data_size = mem.memory_size();
|
||||
device_ptr client_pointer = mem.device_pointer;
|
||||
|
||||
DataVector &data_v = data_vector_find(client_pointer);
|
||||
if(client_pointer) {
|
||||
/* Lookup existing host side data buffer. */
|
||||
DataVector &data_v = data_vector_find(client_pointer);
|
||||
mem.data_pointer = (device_ptr)&data_v[0];
|
||||
|
||||
size_t data_size = mem.memory_size();
|
||||
/* Translate the client pointer to a real device pointer. */
|
||||
mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
|
||||
}
|
||||
else {
|
||||
/* Allocate host side data buffer. */
|
||||
DataVector &data_v = data_vector_insert(client_pointer, data_size);
|
||||
mem.data_pointer = (data_size)? (device_ptr)&(data_v[0]): 0;
|
||||
}
|
||||
|
||||
/* get pointer to memory buffer for device buffer */
|
||||
mem.data_pointer = (device_ptr)&data_v[0];
|
||||
|
||||
/* copy data from network into memory buffer */
|
||||
/* Copy data from network into memory buffer. */
|
||||
rcv.read_buffer((uint8_t*)mem.data_pointer, data_size);
|
||||
|
||||
/* translate the client pointer to a real device pointer */
|
||||
mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
|
||||
|
||||
/* copy the data from the memory buffer to the device buffer */
|
||||
/* Copy the data from the memory buffer to the device buffer. */
|
||||
device->mem_copy_to(mem);
|
||||
|
||||
if(!client_pointer) {
|
||||
/* Store a mapping to/from client_pointer and real device pointer. */
|
||||
pointer_mapping_insert(client_pointer, mem.device_pointer);
|
||||
}
|
||||
}
|
||||
else if(rcv.name == "mem_copy_from") {
|
||||
string name;
|
||||
@ -538,14 +514,30 @@ protected:
|
||||
rcv.read(mem, name);
|
||||
lock.unlock();
|
||||
|
||||
size_t data_size = mem.memory_size();
|
||||
device_ptr client_pointer = mem.device_pointer;
|
||||
mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
|
||||
|
||||
DataVector &data_v = data_vector_find(client_pointer);
|
||||
if(client_pointer) {
|
||||
/* Lookup existing host side data buffer. */
|
||||
DataVector &data_v = data_vector_find(client_pointer);
|
||||
mem.data_pointer = (device_ptr)&data_v[0];
|
||||
|
||||
mem.data_pointer = (device_ptr)&(data_v[0]);
|
||||
/* Translate the client pointer to a real device pointer. */
|
||||
mem.device_pointer = device_ptr_from_client_pointer(client_pointer);
|
||||
}
|
||||
else {
|
||||
/* Allocate host side data buffer. */
|
||||
DataVector &data_v = data_vector_insert(client_pointer, data_size);
|
||||
mem.data_pointer = (data_size)? (device_ptr)&(data_v[0]): 0;
|
||||
}
|
||||
|
||||
/* Zero memory. */
|
||||
device->mem_zero(mem);
|
||||
|
||||
if(!client_pointer) {
|
||||
/* Store a mapping to/from client_pointer and real device pointer. */
|
||||
pointer_mapping_insert(client_pointer, mem.device_pointer);
|
||||
}
|
||||
}
|
||||
else if(rcv.name == "mem_free") {
|
||||
string name;
|
||||
@ -573,45 +565,6 @@ protected:
|
||||
|
||||
device->const_copy_to(name_string.c_str(), &host_vector[0], size);
|
||||
}
|
||||
else if(rcv.name == "tex_alloc") {
|
||||
string name;
|
||||
network_device_memory mem(device);
|
||||
device_ptr client_pointer;
|
||||
|
||||
rcv.read(mem, name);
|
||||
lock.unlock();
|
||||
|
||||
client_pointer = mem.device_pointer;
|
||||
|
||||
size_t data_size = mem.memory_size();
|
||||
|
||||
DataVector &data_v = data_vector_insert(client_pointer, data_size);
|
||||
|
||||
if(data_size)
|
||||
mem.data_pointer = (device_ptr)&(data_v[0]);
|
||||
else
|
||||
mem.data_pointer = 0;
|
||||
|
||||
rcv.read_buffer((uint8_t*)mem.data_pointer, data_size);
|
||||
|
||||
device->tex_alloc(mem);
|
||||
|
||||
pointer_mapping_insert(client_pointer, mem.device_pointer);
|
||||
}
|
||||
else if(rcv.name == "tex_free") {
|
||||
string name;
|
||||
network_device_memory mem(device);
|
||||
device_ptr client_pointer;
|
||||
|
||||
rcv.read(mem, name);
|
||||
lock.unlock();
|
||||
|
||||
client_pointer = mem.device_pointer;
|
||||
|
||||
mem.device_pointer = device_ptr_from_client_pointer_erase(client_pointer);
|
||||
|
||||
device->tex_free(mem);
|
||||
}
|
||||
else if(rcv.name == "load_kernels") {
|
||||
DeviceRequestedFeatures requested_features;
|
||||
rcv.read(requested_features.experimental);
|
||||
@ -696,7 +649,7 @@ protected:
|
||||
}
|
||||
}
|
||||
|
||||
bool task_acquire_tile(Device *device, RenderTile& tile)
|
||||
bool task_acquire_tile(Device *, RenderTile& tile)
|
||||
{
|
||||
thread_scoped_lock acquire_lock(acquire_mutex);
|
||||
|
||||
|
@ -279,6 +279,11 @@ public:
|
||||
|
||||
mem.name = name.c_str();
|
||||
mem.data_pointer = 0;
|
||||
|
||||
/* Can't transfer OpenGL texture over network. */
|
||||
if(mem.type == MEM_PIXELS) {
|
||||
mem.type = MEM_WRITE_ONLY;
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T> void read(T& data)
|
||||
|
@ -61,11 +61,11 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device)
|
||||
|
||||
DeviceSplitKernel::~DeviceSplitKernel()
|
||||
{
|
||||
device->mem_free(split_data);
|
||||
device->mem_free(ray_state);
|
||||
device->mem_free(use_queues_flag);
|
||||
device->mem_free(queue_index);
|
||||
device->mem_free(work_pool_wgs);
|
||||
split_data.free();
|
||||
ray_state.free();
|
||||
use_queues_flag.free();
|
||||
queue_index.free();
|
||||
work_pool_wgs.free();
|
||||
|
||||
delete kernel_path_init;
|
||||
delete kernel_scene_intersect;
|
||||
@ -175,20 +175,11 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
|
||||
unsigned int max_work_groups = num_global_elements / work_pool_size + 1;
|
||||
|
||||
/* Allocate work_pool_wgs memory. */
|
||||
work_pool_wgs.resize(max_work_groups);
|
||||
device->mem_alloc(work_pool_wgs);
|
||||
|
||||
queue_index.resize(NUM_QUEUES);
|
||||
device->mem_alloc(queue_index);
|
||||
|
||||
use_queues_flag.resize(1);
|
||||
device->mem_alloc(use_queues_flag);
|
||||
|
||||
ray_state.resize(num_global_elements);
|
||||
device->mem_alloc(ray_state);
|
||||
|
||||
split_data.resize(state_buffer_size(kgbuffer, kernel_data, num_global_elements));
|
||||
device->mem_alloc(split_data);
|
||||
work_pool_wgs.alloc_to_device(max_work_groups);
|
||||
queue_index.alloc_to_device(NUM_QUEUES);
|
||||
use_queues_flag.alloc_to_device(1);
|
||||
split_data.alloc_to_device(state_buffer_size(kgbuffer, kernel_data, num_global_elements));
|
||||
ray_state.alloc(num_global_elements);
|
||||
}
|
||||
|
||||
#define ENQUEUE_SPLIT_KERNEL(name, global_size, local_size) \
|
||||
@ -225,9 +216,9 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
|
||||
/* reset state memory here as global size for data_init
|
||||
* kernel might not be large enough to do in kernel
|
||||
*/
|
||||
device->mem_zero(work_pool_wgs);
|
||||
device->mem_zero(split_data);
|
||||
device->mem_zero(ray_state);
|
||||
work_pool_wgs.zero_to_device();
|
||||
split_data.zero_to_device();
|
||||
ray_state.zero_to_device();
|
||||
|
||||
if(!enqueue_split_kernel_data_init(KernelDimensions(global_size, local_size),
|
||||
subtile,
|
||||
@ -284,7 +275,7 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
|
||||
}
|
||||
|
||||
/* Decide if we should exit path-iteration in host. */
|
||||
device->mem_copy_from(ray_state, 0, global_size[0] * global_size[1] * sizeof(char), 1, 1);
|
||||
ray_state.copy_from_device(0, global_size[0] * global_size[1], 1);
|
||||
|
||||
activeRaysAvailable = false;
|
||||
|
||||
|
@ -76,8 +76,7 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
|
||||
device_only_memory<uchar> *new_buffer =
|
||||
new device_only_memory<uchar>(device, "memory manager buffer");
|
||||
|
||||
new_buffer->resize(total_size);
|
||||
device->mem_alloc(*new_buffer);
|
||||
new_buffer->alloc_to_device(total_size);
|
||||
|
||||
size_t offset = 0;
|
||||
|
||||
@ -111,7 +110,6 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
|
||||
offset += allocation->size;
|
||||
}
|
||||
|
||||
device->mem_free(*buffer);
|
||||
delete buffer;
|
||||
|
||||
buffer = new_buffer;
|
||||
@ -144,9 +142,9 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
|
||||
clFinish(device->cqCommandQueue);
|
||||
}
|
||||
|
||||
void MemoryManager::DeviceBuffer::free(OpenCLDeviceBase *device)
|
||||
void MemoryManager::DeviceBuffer::free(OpenCLDeviceBase *)
|
||||
{
|
||||
device->mem_free(*buffer);
|
||||
buffer->free();
|
||||
}
|
||||
|
||||
MemoryManager::DeviceBuffer* MemoryManager::smallest_device_buffer()
|
||||
|
@ -74,7 +74,7 @@ void OpenCLDeviceBase::opencl_assert_err(cl_int err, const char* where)
|
||||
OpenCLDeviceBase::OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
|
||||
: Device(info, stats, background_),
|
||||
memory_manager(this),
|
||||
texture_info_buffer(this, "__texture_info", MEM_READ_ONLY)
|
||||
texture_info(this, "__texture_info", MEM_TEXTURE)
|
||||
{
|
||||
cpPlatform = NULL;
|
||||
cdDevice = NULL;
|
||||
@ -157,7 +157,6 @@ OpenCLDeviceBase::~OpenCLDeviceBase()
|
||||
|
||||
ConstMemMap::iterator mt;
|
||||
for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
|
||||
mem_free(*(mt->second));
|
||||
delete mt->second;
|
||||
}
|
||||
|
||||
@ -318,9 +317,9 @@ void OpenCLDeviceBase::mem_alloc(device_memory& mem)
|
||||
cl_mem_flags mem_flag;
|
||||
void *mem_ptr = NULL;
|
||||
|
||||
if(mem.type == MEM_READ_ONLY)
|
||||
if(mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE)
|
||||
mem_flag = CL_MEM_READ_ONLY;
|
||||
else if(mem.type == MEM_WRITE_ONLY)
|
||||
else if(mem.type == MEM_WRITE_ONLY || mem.type == MEM_PIXELS)
|
||||
mem_flag = CL_MEM_WRITE_ONLY;
|
||||
else
|
||||
mem_flag = CL_MEM_READ_WRITE;
|
||||
@ -348,17 +347,27 @@ void OpenCLDeviceBase::mem_alloc(device_memory& mem)
|
||||
|
||||
void OpenCLDeviceBase::mem_copy_to(device_memory& mem)
|
||||
{
|
||||
/* this is blocking */
|
||||
size_t size = mem.memory_size();
|
||||
if(size != 0) {
|
||||
opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
|
||||
CL_MEM_PTR(mem.device_pointer),
|
||||
CL_TRUE,
|
||||
0,
|
||||
size,
|
||||
(void*)mem.data_pointer,
|
||||
0,
|
||||
NULL, NULL));
|
||||
if(mem.type == MEM_TEXTURE) {
|
||||
tex_free(mem);
|
||||
tex_alloc(mem);
|
||||
}
|
||||
else {
|
||||
if(!mem.device_pointer) {
|
||||
mem_alloc(mem);
|
||||
}
|
||||
|
||||
/* this is blocking */
|
||||
size_t size = mem.memory_size();
|
||||
if(size != 0) {
|
||||
opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
|
||||
CL_MEM_PTR(mem.device_pointer),
|
||||
CL_TRUE,
|
||||
0,
|
||||
size,
|
||||
(void*)mem.data_pointer,
|
||||
0,
|
||||
NULL, NULL));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -410,6 +419,10 @@ void OpenCLDeviceBase::mem_zero_kernel(device_ptr mem, size_t size)
|
||||
|
||||
void OpenCLDeviceBase::mem_zero(device_memory& mem)
|
||||
{
|
||||
if(!mem.device_pointer) {
|
||||
mem_alloc(mem);
|
||||
}
|
||||
|
||||
if(mem.device_pointer) {
|
||||
if(base_program.is_loaded()) {
|
||||
mem_zero_kernel(mem.device_pointer, mem.memory_size());
|
||||
@ -445,14 +458,19 @@ void OpenCLDeviceBase::mem_zero(device_memory& mem)
|
||||
|
||||
void OpenCLDeviceBase::mem_free(device_memory& mem)
|
||||
{
|
||||
if(mem.device_pointer) {
|
||||
if(mem.device_pointer != null_mem) {
|
||||
opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
|
||||
}
|
||||
mem.device_pointer = 0;
|
||||
if(mem.type == MEM_TEXTURE) {
|
||||
tex_free(mem);
|
||||
}
|
||||
else {
|
||||
if(mem.device_pointer) {
|
||||
if(mem.device_pointer != null_mem) {
|
||||
opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
|
||||
}
|
||||
mem.device_pointer = 0;
|
||||
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_size = 0;
|
||||
stats.mem_free(mem.device_size);
|
||||
mem.device_size = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -464,9 +482,9 @@ int OpenCLDeviceBase::mem_address_alignment()
|
||||
device_ptr OpenCLDeviceBase::mem_alloc_sub_ptr(device_memory& mem, int offset, int size)
|
||||
{
|
||||
cl_mem_flags mem_flag;
|
||||
if(mem.type == MEM_READ_ONLY)
|
||||
if(mem.type == MEM_READ_ONLY || mem.type == MEM_TEXTURE)
|
||||
mem_flag = CL_MEM_READ_ONLY;
|
||||
else if(mem.type == MEM_WRITE_ONLY)
|
||||
else if(mem.type == MEM_WRITE_ONLY || mem.type == MEM_PIXELS)
|
||||
mem_flag = CL_MEM_WRITE_ONLY;
|
||||
else
|
||||
mem_flag = CL_MEM_READ_WRITE;
|
||||
@ -498,9 +516,7 @@ void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
|
||||
|
||||
if(i == const_mem_map.end()) {
|
||||
data = new device_vector<uchar>(this, name, MEM_READ_ONLY);
|
||||
data->resize(size);
|
||||
|
||||
mem_alloc(*data);
|
||||
data->alloc(size);
|
||||
const_mem_map.insert(ConstMemMap::value_type(name, data));
|
||||
}
|
||||
else {
|
||||
@ -508,7 +524,7 @@ void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
|
||||
}
|
||||
|
||||
memcpy(data->get_data(), host, size);
|
||||
mem_copy_to(*data);
|
||||
data->copy_to_device();
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::tex_alloc(device_memory& mem)
|
||||
@ -1037,8 +1053,7 @@ bool OpenCLDeviceBase::denoising_detect_outliers(device_ptr image_ptr,
|
||||
bool OpenCLDeviceBase::denoising_set_tiles(device_ptr *buffers,
|
||||
DenoisingTask *task)
|
||||
{
|
||||
mem_alloc(task->tiles_mem);
|
||||
mem_copy_to(task->tiles_mem);
|
||||
task->tiles_mem.copy_to_device();
|
||||
|
||||
cl_mem tiles_mem = CL_MEM_PTR(task->tiles_mem.device_pointer);
|
||||
|
||||
|
@ -128,8 +128,7 @@ public:
|
||||
|
||||
/* Allocate buffer for kernel globals */
|
||||
device_only_memory<KernelGlobals> kgbuffer(this, "kernel_globals");
|
||||
kgbuffer.resize(1);
|
||||
mem_alloc(kgbuffer);
|
||||
kgbuffer.alloc_to_device(1);
|
||||
|
||||
/* Keep rendering tiles until done. */
|
||||
while(task->acquire_tile(this, tile)) {
|
||||
@ -160,7 +159,7 @@ public:
|
||||
task->release_tile(tile);
|
||||
}
|
||||
|
||||
mem_free(kgbuffer);
|
||||
kgbuffer.free();
|
||||
}
|
||||
}
|
||||
|
||||
@ -289,8 +288,8 @@ public:
|
||||
virtual uint64_t state_buffer_size(device_memory& kg, device_memory& data, size_t num_threads)
|
||||
{
|
||||
device_vector<uint64_t> size_buffer(device, "size_buffer", MEM_READ_WRITE);
|
||||
size_buffer.resize(1);
|
||||
device->mem_alloc(size_buffer);
|
||||
size_buffer.alloc(1);
|
||||
size_buffer.zero_to_device();
|
||||
|
||||
uint threads = num_threads;
|
||||
device->kernel_set_args(device->program_state_buffer_size(), 0, kg, data, threads, size_buffer);
|
||||
@ -308,9 +307,9 @@ public:
|
||||
|
||||
device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel");
|
||||
|
||||
device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t));
|
||||
size_buffer.copy_from_device(0, 1, 1);
|
||||
size_t size = size_buffer[0];
|
||||
device->mem_free(size_buffer);
|
||||
size_buffer.free();
|
||||
|
||||
if(device->ciErr != CL_SUCCESS) {
|
||||
string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()",
|
||||
|
@ -151,7 +151,7 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
|
||||
|
||||
/* setup input for device task */
|
||||
device_vector<uint4> d_input(device, "bake_input", MEM_READ_ONLY);
|
||||
uint4 *d_input_data = d_input.resize(shader_size * 2);
|
||||
uint4 *d_input_data = d_input.alloc(shader_size * 2);
|
||||
size_t d_input_size = 0;
|
||||
|
||||
for(size_t i = shader_offset; i < (shader_offset + shader_size); i++) {
|
||||
@ -166,16 +166,13 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
|
||||
|
||||
/* run device task */
|
||||
device_vector<float4> d_output(device, "bake_output", MEM_READ_WRITE);
|
||||
d_output.resize(shader_size);
|
||||
d_output.alloc(shader_size);
|
||||
d_output.zero_to_device();
|
||||
d_input.copy_to_device();
|
||||
|
||||
/* needs to be up to data for attribute access */
|
||||
device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
|
||||
|
||||
device->mem_alloc(d_input);
|
||||
device->mem_copy_to(d_input);
|
||||
device->mem_alloc(d_output);
|
||||
device->mem_zero(d_output);
|
||||
|
||||
DeviceTask task(DeviceTask::SHADER);
|
||||
task.shader_input = d_input.device_pointer;
|
||||
task.shader_output = d_output.device_pointer;
|
||||
@ -192,15 +189,14 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
|
||||
device->task_wait();
|
||||
|
||||
if(progress.get_cancel()) {
|
||||
device->mem_free(d_input);
|
||||
device->mem_free(d_output);
|
||||
d_input.free();
|
||||
d_output.free();
|
||||
m_is_baking = false;
|
||||
return false;
|
||||
}
|
||||
|
||||
device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float4));
|
||||
device->mem_free(d_input);
|
||||
device->mem_free(d_output);
|
||||
d_output.copy_from_device(0, 1, d_output.size());
|
||||
d_input.free();
|
||||
|
||||
/* read result */
|
||||
int k = 0;
|
||||
@ -218,6 +214,8 @@ bool BakeManager::bake(Device *device, DeviceScene *dscene, Scene *scene, Progre
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
d_output.free();
|
||||
}
|
||||
|
||||
m_is_baking = false;
|
||||
|
@ -115,54 +115,35 @@ RenderTile::RenderTile()
|
||||
/* Render Buffers */
|
||||
|
||||
RenderBuffers::RenderBuffers(Device *device)
|
||||
: buffer(device, "RenderBuffers", MEM_READ_WRITE),
|
||||
device(device)
|
||||
: buffer(device, "RenderBuffers", MEM_READ_WRITE)
|
||||
{
|
||||
}
|
||||
|
||||
RenderBuffers::~RenderBuffers()
|
||||
{
|
||||
device_free();
|
||||
buffer.free();
|
||||
}
|
||||
|
||||
void RenderBuffers::device_free()
|
||||
{
|
||||
if(buffer.device_pointer) {
|
||||
device->mem_free(buffer);
|
||||
buffer.clear();
|
||||
}
|
||||
}
|
||||
|
||||
void RenderBuffers::reset(Device *device, BufferParams& params_)
|
||||
void RenderBuffers::reset(BufferParams& params_)
|
||||
{
|
||||
params = params_;
|
||||
|
||||
/* free existing buffers */
|
||||
device_free();
|
||||
|
||||
/* allocate buffer */
|
||||
buffer.resize(params.width*params.height*params.get_passes_size());
|
||||
device->mem_alloc(buffer);
|
||||
device->mem_zero(buffer);
|
||||
/* re-allocate buffer */
|
||||
buffer.alloc(params.width*params.height*params.get_passes_size());
|
||||
buffer.zero_to_device();
|
||||
}
|
||||
|
||||
void RenderBuffers::zero(Device *device)
|
||||
void RenderBuffers::zero()
|
||||
{
|
||||
if(buffer.device_pointer) {
|
||||
device->mem_zero(buffer);
|
||||
}
|
||||
buffer.zero_to_device();
|
||||
}
|
||||
|
||||
bool RenderBuffers::copy_from_device(Device *from_device)
|
||||
bool RenderBuffers::copy_from_device()
|
||||
{
|
||||
if(!buffer.device_pointer)
|
||||
return false;
|
||||
|
||||
if(!from_device) {
|
||||
from_device = device;
|
||||
}
|
||||
|
||||
from_device->mem_copy_from(buffer, 0, params.width, params.height, params.get_passes_size()*sizeof(float));
|
||||
buffer.copy_from_device(0, params.width * params.get_passes_size(), params.height);
|
||||
|
||||
return true;
|
||||
}
|
||||
@ -402,47 +383,30 @@ DisplayBuffer::DisplayBuffer(Device *device, bool linear)
|
||||
draw_height(0),
|
||||
transparent(true), /* todo: determine from background */
|
||||
half_float(linear),
|
||||
rgba_byte(device, "display buffer byte", MEM_WRITE_ONLY),
|
||||
rgba_half(device, "display buffer half", MEM_WRITE_ONLY),
|
||||
device(device)
|
||||
rgba_byte(device, "display buffer byte"),
|
||||
rgba_half(device, "display buffer half")
|
||||
{
|
||||
}
|
||||
|
||||
DisplayBuffer::~DisplayBuffer()
|
||||
{
|
||||
device_free();
|
||||
rgba_byte.free();
|
||||
rgba_half.free();
|
||||
}
|
||||
|
||||
void DisplayBuffer::device_free()
|
||||
{
|
||||
if(rgba_byte.device_pointer) {
|
||||
device->pixels_free(rgba_byte);
|
||||
rgba_byte.clear();
|
||||
}
|
||||
if(rgba_half.device_pointer) {
|
||||
device->pixels_free(rgba_half);
|
||||
rgba_half.clear();
|
||||
}
|
||||
}
|
||||
|
||||
void DisplayBuffer::reset(Device *device, BufferParams& params_)
|
||||
void DisplayBuffer::reset(BufferParams& params_)
|
||||
{
|
||||
draw_width = 0;
|
||||
draw_height = 0;
|
||||
|
||||
params = params_;
|
||||
|
||||
/* free existing buffers */
|
||||
device_free();
|
||||
|
||||
/* allocate display pixels */
|
||||
if(half_float) {
|
||||
rgba_half.resize(params.width, params.height);
|
||||
device->pixels_alloc(rgba_half);
|
||||
rgba_half.alloc_to_device(params.width, params.height);
|
||||
}
|
||||
else {
|
||||
rgba_byte.resize(params.width, params.height);
|
||||
device->pixels_alloc(rgba_byte);
|
||||
rgba_byte.alloc_to_device(params.width, params.height);
|
||||
}
|
||||
}
|
||||
|
||||
@ -457,7 +421,8 @@ void DisplayBuffer::draw_set(int width, int height)
|
||||
void DisplayBuffer::draw(Device *device, const DeviceDrawParams& draw_params)
|
||||
{
|
||||
if(draw_width != 0 && draw_height != 0) {
|
||||
device_memory& rgba = rgba_data();
|
||||
device_memory& rgba = (half_float)? (device_memory&)rgba_half:
|
||||
(device_memory&)rgba_byte;
|
||||
|
||||
device->draw_pixels(rgba, 0, draw_width, draw_height, params.full_x, params.full_y, params.width, params.height, transparent, draw_params);
|
||||
}
|
||||
@ -468,7 +433,7 @@ bool DisplayBuffer::draw_ready()
|
||||
return (draw_width != 0 && draw_height != 0);
|
||||
}
|
||||
|
||||
void DisplayBuffer::write(Device *device, const string& filename)
|
||||
void DisplayBuffer::write(const string& filename)
|
||||
{
|
||||
int w = draw_width;
|
||||
int h = draw_height;
|
||||
@ -480,21 +445,19 @@ void DisplayBuffer::write(Device *device, const string& filename)
|
||||
return;
|
||||
|
||||
/* read buffer from device */
|
||||
device_memory& rgba = rgba_data();
|
||||
device->pixels_copy_from(rgba, 0, w, h);
|
||||
uchar4 *pixels = rgba_byte.copy_from_device(0, w, h);
|
||||
|
||||
/* write image */
|
||||
ImageOutput *out = ImageOutput::create(filename);
|
||||
ImageSpec spec(w, h, 4, TypeDesc::UINT8);
|
||||
int scanlinesize = w*4*sizeof(uchar);
|
||||
|
||||
out->open(filename, spec);
|
||||
|
||||
/* conversion for different top/bottom convention */
|
||||
out->write_image(TypeDesc::UINT8,
|
||||
(uchar*)rgba.data_pointer + (h-1)*scanlinesize,
|
||||
(uchar*)(pixels + (h-1)*w),
|
||||
AutoStride,
|
||||
-scanlinesize,
|
||||
-w*sizeof(uchar4),
|
||||
AutoStride);
|
||||
|
||||
out->close();
|
||||
@ -502,13 +465,5 @@ void DisplayBuffer::write(Device *device, const string& filename)
|
||||
delete out;
|
||||
}
|
||||
|
||||
device_memory& DisplayBuffer::rgba_data()
|
||||
{
|
||||
if(half_float)
|
||||
return rgba_half;
|
||||
else
|
||||
return rgba_byte;
|
||||
}
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
||||
|
@ -75,20 +75,15 @@ public:
|
||||
/* float buffer */
|
||||
device_vector<float> buffer;
|
||||
|
||||
Device *device;
|
||||
|
||||
explicit RenderBuffers(Device *device);
|
||||
~RenderBuffers();
|
||||
|
||||
void reset(Device *device, BufferParams& params);
|
||||
void zero(Device *device);
|
||||
void reset(BufferParams& params);
|
||||
void zero();
|
||||
|
||||
bool copy_from_device(Device *from_device = NULL);
|
||||
bool copy_from_device();
|
||||
bool get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels);
|
||||
bool get_denoising_pass_rect(int offset, float exposure, int sample, int components, float *pixels);
|
||||
|
||||
protected:
|
||||
void device_free();
|
||||
};
|
||||
|
||||
/* Display Buffer
|
||||
@ -109,25 +104,18 @@ public:
|
||||
/* use half float? */
|
||||
bool half_float;
|
||||
/* byte buffer for converted result */
|
||||
device_vector<uchar4> rgba_byte;
|
||||
device_vector<half4> rgba_half;
|
||||
device_pixels<uchar4> rgba_byte;
|
||||
device_pixels<half4> rgba_half;
|
||||
|
||||
DisplayBuffer(Device *device, bool linear = false);
|
||||
~DisplayBuffer();
|
||||
|
||||
void reset(Device *device, BufferParams& params);
|
||||
void write(Device *device, const string& filename);
|
||||
void reset(BufferParams& params);
|
||||
void write(const string& filename);
|
||||
|
||||
void draw_set(int width, int height);
|
||||
void draw(Device *device, const DeviceDrawParams& draw_params);
|
||||
bool draw_ready();
|
||||
|
||||
device_memory& rgba_data();
|
||||
|
||||
protected:
|
||||
void device_free();
|
||||
|
||||
Device *device;
|
||||
};
|
||||
|
||||
/* Render Tile
|
||||
|
@ -532,7 +532,8 @@ bool ImageManager::file_load_image(Image *img,
|
||||
pixels = &pixels_storage[0];
|
||||
}
|
||||
else {
|
||||
pixels = (StorageType*)tex_img.resize(width, height, depth);
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
pixels = (StorageType*)tex_img.alloc(width, height, depth);
|
||||
}
|
||||
if(pixels == NULL) {
|
||||
/* Could be that we've run out of memory. */
|
||||
@ -686,9 +687,16 @@ bool ImageManager::file_load_image(Image *img,
|
||||
scale_factor,
|
||||
&scaled_pixels,
|
||||
&scaled_width, &scaled_height, &scaled_depth);
|
||||
StorageType *texture_pixels = (StorageType*)tex_img.resize(scaled_width,
|
||||
scaled_height,
|
||||
scaled_depth);
|
||||
|
||||
StorageType *texture_pixels;
|
||||
|
||||
{
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
texture_pixels = (StorageType*)tex_img.alloc(scaled_width,
|
||||
scaled_height,
|
||||
scaled_depth);
|
||||
}
|
||||
|
||||
memcpy(texture_pixels,
|
||||
&scaled_pixels[0],
|
||||
scaled_pixels.size() * sizeof(StorageType));
|
||||
@ -722,14 +730,14 @@ void ImageManager::device_load_image(Device *device,
|
||||
/* Free previous texture in slot. */
|
||||
if(img->mem) {
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
device->tex_free(*img->mem);
|
||||
delete img->mem;
|
||||
img->mem = NULL;
|
||||
}
|
||||
|
||||
/* Create new texture. */
|
||||
if(type == IMAGE_DATA_TYPE_FLOAT4) {
|
||||
device_vector<float4> *tex_img = new device_vector<float4>(device, name.c_str());
|
||||
device_vector<float4> *tex_img
|
||||
= new device_vector<float4>(device, name.c_str(), MEM_TEXTURE);
|
||||
|
||||
if(!file_load_image<TypeDesc::FLOAT, float>(img,
|
||||
type,
|
||||
@ -737,7 +745,7 @@ void ImageManager::device_load_image(Device *device,
|
||||
*tex_img))
|
||||
{
|
||||
/* on failure to load, we set a 1x1 pixels pink image */
|
||||
float *pixels = (float*)tex_img->resize(1, 1);
|
||||
float *pixels = (float*)tex_img->alloc(1, 1);
|
||||
|
||||
pixels[0] = TEX_IMAGE_MISSING_R;
|
||||
pixels[1] = TEX_IMAGE_MISSING_G;
|
||||
@ -746,9 +754,15 @@ void ImageManager::device_load_image(Device *device,
|
||||
}
|
||||
|
||||
img->mem = tex_img;
|
||||
img->mem->interpolation = img->interpolation;
|
||||
img->mem->extension = img->extension;
|
||||
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
tex_img->copy_to_device();
|
||||
}
|
||||
else if(type == IMAGE_DATA_TYPE_FLOAT) {
|
||||
device_vector<float> *tex_img = new device_vector<float>(device, name.c_str());
|
||||
device_vector<float> *tex_img
|
||||
= new device_vector<float>(device, name.c_str(), MEM_TEXTURE);
|
||||
|
||||
if(!file_load_image<TypeDesc::FLOAT, float>(img,
|
||||
type,
|
||||
@ -756,15 +770,21 @@ void ImageManager::device_load_image(Device *device,
|
||||
*tex_img))
|
||||
{
|
||||
/* on failure to load, we set a 1x1 pixels pink image */
|
||||
float *pixels = (float*)tex_img->resize(1, 1);
|
||||
float *pixels = (float*)tex_img->alloc(1, 1);
|
||||
|
||||
pixels[0] = TEX_IMAGE_MISSING_R;
|
||||
}
|
||||
|
||||
img->mem = tex_img;
|
||||
img->mem->interpolation = img->interpolation;
|
||||
img->mem->extension = img->extension;
|
||||
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
tex_img->copy_to_device();
|
||||
}
|
||||
else if(type == IMAGE_DATA_TYPE_BYTE4) {
|
||||
device_vector<uchar4> *tex_img = new device_vector<uchar4>(device, name.c_str());
|
||||
device_vector<uchar4> *tex_img
|
||||
= new device_vector<uchar4>(device, name.c_str(), MEM_TEXTURE);
|
||||
|
||||
if(!file_load_image<TypeDesc::UINT8, uchar>(img,
|
||||
type,
|
||||
@ -772,7 +792,7 @@ void ImageManager::device_load_image(Device *device,
|
||||
*tex_img))
|
||||
{
|
||||
/* on failure to load, we set a 1x1 pixels pink image */
|
||||
uchar *pixels = (uchar*)tex_img->resize(1, 1);
|
||||
uchar *pixels = (uchar*)tex_img->alloc(1, 1);
|
||||
|
||||
pixels[0] = (TEX_IMAGE_MISSING_R * 255);
|
||||
pixels[1] = (TEX_IMAGE_MISSING_G * 255);
|
||||
@ -781,31 +801,43 @@ void ImageManager::device_load_image(Device *device,
|
||||
}
|
||||
|
||||
img->mem = tex_img;
|
||||
img->mem->interpolation = img->interpolation;
|
||||
img->mem->extension = img->extension;
|
||||
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
tex_img->copy_to_device();
|
||||
}
|
||||
else if(type == IMAGE_DATA_TYPE_BYTE) {
|
||||
device_vector<uchar> *tex_img = new device_vector<uchar>(device, name.c_str());
|
||||
device_vector<uchar> *tex_img
|
||||
= new device_vector<uchar>(device, name.c_str(), MEM_TEXTURE);
|
||||
|
||||
if(!file_load_image<TypeDesc::UINT8, uchar>(img,
|
||||
type,
|
||||
texture_limit,
|
||||
*tex_img)) {
|
||||
/* on failure to load, we set a 1x1 pixels pink image */
|
||||
uchar *pixels = (uchar*)tex_img->resize(1, 1);
|
||||
uchar *pixels = (uchar*)tex_img->alloc(1, 1);
|
||||
|
||||
pixels[0] = (TEX_IMAGE_MISSING_R * 255);
|
||||
}
|
||||
|
||||
img->mem = tex_img;
|
||||
img->mem->interpolation = img->interpolation;
|
||||
img->mem->extension = img->extension;
|
||||
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
tex_img->copy_to_device();
|
||||
}
|
||||
else if(type == IMAGE_DATA_TYPE_HALF4) {
|
||||
device_vector<half4> *tex_img = new device_vector<half4>(device, name.c_str());
|
||||
device_vector<half4> *tex_img
|
||||
= new device_vector<half4>(device, name.c_str(), MEM_TEXTURE);
|
||||
|
||||
if(!file_load_image<TypeDesc::HALF, half>(img,
|
||||
type,
|
||||
texture_limit,
|
||||
*tex_img)) {
|
||||
/* on failure to load, we set a 1x1 pixels pink image */
|
||||
half *pixels = (half*)tex_img->resize(1, 1);
|
||||
half *pixels = (half*)tex_img->alloc(1, 1);
|
||||
|
||||
pixels[0] = TEX_IMAGE_MISSING_R;
|
||||
pixels[1] = TEX_IMAGE_MISSING_G;
|
||||
@ -814,37 +846,38 @@ void ImageManager::device_load_image(Device *device,
|
||||
}
|
||||
|
||||
img->mem = tex_img;
|
||||
img->mem->interpolation = img->interpolation;
|
||||
img->mem->extension = img->extension;
|
||||
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
tex_img->copy_to_device();
|
||||
}
|
||||
else if(type == IMAGE_DATA_TYPE_HALF) {
|
||||
device_vector<half> *tex_img = new device_vector<half>(device, name.c_str());
|
||||
device_vector<half> *tex_img
|
||||
= new device_vector<half>(device, name.c_str(), MEM_TEXTURE);
|
||||
|
||||
if(!file_load_image<TypeDesc::HALF, half>(img,
|
||||
type,
|
||||
texture_limit,
|
||||
*tex_img)) {
|
||||
/* on failure to load, we set a 1x1 pixels pink image */
|
||||
half *pixels = (half*)tex_img->resize(1, 1);
|
||||
half *pixels = (half*)tex_img->alloc(1, 1);
|
||||
|
||||
pixels[0] = TEX_IMAGE_MISSING_R;
|
||||
}
|
||||
|
||||
img->mem = tex_img;
|
||||
}
|
||||
|
||||
/* Copy to device. */
|
||||
if(img->mem) {
|
||||
img->mem->interpolation = img->interpolation;
|
||||
img->mem->extension = img->extension;
|
||||
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
device->tex_alloc(*img->mem);
|
||||
tex_img->copy_to_device();
|
||||
}
|
||||
|
||||
|
||||
img->need_load = false;
|
||||
}
|
||||
|
||||
void ImageManager::device_free_image(Device *device, ImageDataType type, int slot)
|
||||
void ImageManager::device_free_image(Device *, ImageDataType type, int slot)
|
||||
{
|
||||
Image *img = images[type][slot];
|
||||
|
||||
@ -858,7 +891,6 @@ void ImageManager::device_free_image(Device *device, ImageDataType type, int slo
|
||||
|
||||
if(img->mem) {
|
||||
thread_scoped_lock device_lock(device_mutex);
|
||||
device->tex_free(*img->mem);
|
||||
delete img->mem;
|
||||
}
|
||||
|
||||
|
@ -191,11 +191,11 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
|
||||
int dimensions = PRNG_BASE_NUM + max_samples*PRNG_BOUNCE_NUM;
|
||||
dimensions = min(dimensions, SOBOL_MAX_DIMENSIONS);
|
||||
|
||||
uint *directions = dscene->sobol_directions.resize(SOBOL_BITS*dimensions);
|
||||
uint *directions = dscene->sobol_directions.alloc(SOBOL_BITS*dimensions);
|
||||
|
||||
sobol_generate_direction_vectors((uint(*)[SOBOL_BITS])directions, dimensions);
|
||||
|
||||
device->tex_alloc(dscene->sobol_directions);
|
||||
dscene->sobol_directions.copy_to_device();
|
||||
|
||||
/* Clamping. */
|
||||
bool use_sample_clamp = (sample_clamp_direct != 0.0f ||
|
||||
@ -208,10 +208,9 @@ void Integrator::device_update(Device *device, DeviceScene *dscene, Scene *scene
|
||||
need_update = false;
|
||||
}
|
||||
|
||||
void Integrator::device_free(Device *device, DeviceScene *dscene)
|
||||
void Integrator::device_free(Device *, DeviceScene *dscene)
|
||||
{
|
||||
device->tex_free(dscene->sobol_directions);
|
||||
dscene->sobol_directions.clear();
|
||||
dscene->sobol_directions.free();
|
||||
}
|
||||
|
||||
bool Integrator::modified(const Integrator& integrator)
|
||||
|
@ -39,7 +39,7 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
|
||||
device_vector<uint4> d_input(device, "background_input", MEM_READ_ONLY);
|
||||
device_vector<float4> d_output(device, "background_output", MEM_WRITE_ONLY);
|
||||
|
||||
uint4 *d_input_data = d_input.resize(width*height);
|
||||
uint4 *d_input_data = d_input.alloc(width*height);
|
||||
|
||||
for(int y = 0; y < height; y++) {
|
||||
for(int x = 0; x < width; x++) {
|
||||
@ -52,16 +52,12 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
|
||||
}
|
||||
|
||||
/* compute on device */
|
||||
d_output.resize(width*height);
|
||||
memset((void*)d_output.data_pointer, 0, d_output.memory_size());
|
||||
d_output.alloc(width*height);
|
||||
d_output.zero_to_device();
|
||||
d_input.copy_to_device();
|
||||
|
||||
device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
|
||||
|
||||
device->mem_alloc(d_input);
|
||||
device->mem_copy_to(d_input);
|
||||
device->mem_alloc(d_output);
|
||||
device->mem_zero(d_output);
|
||||
|
||||
DeviceTask main_task(DeviceTask::SHADER);
|
||||
main_task.shader_input = d_input.device_pointer;
|
||||
main_task.shader_output = d_output.device_pointer;
|
||||
@ -78,13 +74,10 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
|
||||
foreach(DeviceTask& task, split_tasks) {
|
||||
device->task_add(task);
|
||||
device->task_wait();
|
||||
device->mem_copy_from(d_output, task.shader_x, 1, task.shader_w, sizeof(float4));
|
||||
d_output.copy_from_device(task.shader_x, 1, task.shader_w);
|
||||
}
|
||||
|
||||
device->mem_free(d_input);
|
||||
device->mem_free(d_output);
|
||||
|
||||
d_input.clear();
|
||||
d_input.free();
|
||||
|
||||
float4 *d_output_data = reinterpret_cast<float4*>(d_output.data_pointer);
|
||||
|
||||
@ -97,6 +90,8 @@ static void shade_background_pixels(Device *device, DeviceScene *dscene, int res
|
||||
pixels[y*width + x].z = d_output_data[y*width + x].z;
|
||||
}
|
||||
}
|
||||
|
||||
d_output.free();
|
||||
}
|
||||
|
||||
/* Light */
|
||||
@ -246,7 +241,7 @@ bool LightManager::object_usable_as_light(Object *object) {
|
||||
return false;
|
||||
}
|
||||
|
||||
void LightManager::device_update_distribution(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
|
||||
void LightManager::device_update_distribution(Device *, DeviceScene *dscene, Scene *scene, Progress& progress)
|
||||
{
|
||||
progress.set_status("Updating Lights", "Computing distribution");
|
||||
|
||||
@ -292,7 +287,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
|
||||
VLOG(1) << "Total " << num_distribution << " of light distribution primitives.";
|
||||
|
||||
/* emission area */
|
||||
float4 *distribution = dscene->light_distribution.resize(num_distribution + 1);
|
||||
float4 *distribution = dscene->light_distribution.alloc(num_distribution + 1);
|
||||
float totarea = 0.0f;
|
||||
|
||||
/* triangles */
|
||||
@ -451,7 +446,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
|
||||
kfilm->pass_shadow_scale *= (float)(num_lights - num_background_lights)/(float)num_lights;
|
||||
|
||||
/* CDF */
|
||||
device->tex_alloc(dscene->light_distribution);
|
||||
dscene->light_distribution.copy_to_device();
|
||||
|
||||
/* Portals */
|
||||
if(num_portals > 0) {
|
||||
@ -466,7 +461,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
|
||||
}
|
||||
}
|
||||
else {
|
||||
dscene->light_distribution.clear();
|
||||
dscene->light_distribution.free();
|
||||
|
||||
kintegrator->num_distribution = 0;
|
||||
kintegrator->num_all_lights = 0;
|
||||
@ -561,8 +556,8 @@ void LightManager::device_update_background(Device *device,
|
||||
|
||||
/* build row distributions and column distribution for the infinite area environment light */
|
||||
int cdf_count = res + 1;
|
||||
float2 *marg_cdf = dscene->light_background_marginal_cdf.resize(cdf_count);
|
||||
float2 *cond_cdf = dscene->light_background_conditional_cdf.resize(cdf_count * cdf_count);
|
||||
float2 *marg_cdf = dscene->light_background_marginal_cdf.alloc(cdf_count);
|
||||
float2 *cond_cdf = dscene->light_background_conditional_cdf.alloc(cdf_count * cdf_count);
|
||||
|
||||
double time_start = time_dt();
|
||||
if(res < 512) {
|
||||
@ -611,11 +606,11 @@ void LightManager::device_update_background(Device *device,
|
||||
VLOG(2) << "Background MIS build time " << time_dt() - time_start << "\n";
|
||||
|
||||
/* update device */
|
||||
device->tex_alloc(dscene->light_background_marginal_cdf);
|
||||
device->tex_alloc(dscene->light_background_conditional_cdf);
|
||||
dscene->light_background_marginal_cdf.copy_to_device();
|
||||
dscene->light_background_conditional_cdf.copy_to_device();
|
||||
}
|
||||
|
||||
void LightManager::device_update_points(Device *device,
|
||||
void LightManager::device_update_points(Device *,
|
||||
DeviceScene *dscene,
|
||||
Scene *scene)
|
||||
{
|
||||
@ -628,7 +623,7 @@ void LightManager::device_update_points(Device *device,
|
||||
}
|
||||
}
|
||||
|
||||
float4 *light_data = dscene->light_data.resize(num_lights*LIGHT_SIZE);
|
||||
float4 *light_data = dscene->light_data.alloc(num_lights*LIGHT_SIZE);
|
||||
|
||||
if(num_lights == 0) {
|
||||
VLOG(1) << "No effective light, ignoring points update.";
|
||||
@ -813,7 +808,7 @@ void LightManager::device_update_points(Device *device,
|
||||
VLOG(1) << "Number of lights without contribution: "
|
||||
<< num_scene_lights - light_index;
|
||||
|
||||
device->tex_alloc(dscene->light_data);
|
||||
dscene->light_data.copy_to_device();
|
||||
}
|
||||
|
||||
void LightManager::device_update(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
|
||||
@ -846,17 +841,12 @@ void LightManager::device_update(Device *device, DeviceScene *dscene, Scene *sce
|
||||
need_update = false;
|
||||
}
|
||||
|
||||
void LightManager::device_free(Device *device, DeviceScene *dscene)
|
||||
void LightManager::device_free(Device *, DeviceScene *dscene)
|
||||
{
|
||||
device->tex_free(dscene->light_distribution);
|
||||
device->tex_free(dscene->light_data);
|
||||
device->tex_free(dscene->light_background_marginal_cdf);
|
||||
device->tex_free(dscene->light_background_conditional_cdf);
|
||||
|
||||
dscene->light_distribution.clear();
|
||||
dscene->light_data.clear();
|
||||
dscene->light_background_marginal_cdf.clear();
|
||||
dscene->light_background_conditional_cdf.clear();
|
||||
dscene->light_distribution.free();
|
||||
dscene->light_data.free();
|
||||
dscene->light_background_marginal_cdf.free();
|
||||
dscene->light_background_conditional_cdf.free();
|
||||
}
|
||||
|
||||
void LightManager::tag_update(Scene * /*scene*/)
|
||||
|
@ -1252,7 +1252,7 @@ void MeshManager::update_osl_attributes(Device *device, Scene *scene, vector<Att
|
||||
#endif
|
||||
}
|
||||
|
||||
void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Scene *scene, vector<AttributeRequestSet>& mesh_attributes)
|
||||
void MeshManager::update_svm_attributes(Device *, DeviceScene *dscene, Scene *scene, vector<AttributeRequestSet>& mesh_attributes)
|
||||
{
|
||||
/* for SVM, the attributes_map table is used to lookup the offset of an
|
||||
* attribute, based on a unique shader attribute id. */
|
||||
@ -1267,7 +1267,7 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
|
||||
return;
|
||||
|
||||
/* create attribute map */
|
||||
uint4 *attr_map = dscene->attributes_map.resize(attr_map_stride*scene->objects.size());
|
||||
uint4 *attr_map = dscene->attributes_map.alloc(attr_map_stride*scene->objects.size());
|
||||
memset(attr_map, 0, dscene->attributes_map.size()*sizeof(uint));
|
||||
|
||||
for(size_t i = 0; i < scene->objects.size(); i++) {
|
||||
@ -1359,7 +1359,7 @@ void MeshManager::update_svm_attributes(Device *device, DeviceScene *dscene, Sce
|
||||
|
||||
/* copy to device */
|
||||
dscene->data.bvh.attributes_map_stride = attr_map_stride;
|
||||
device->tex_alloc(dscene->attributes_map);
|
||||
dscene->attributes_map.copy_to_device();
|
||||
}
|
||||
|
||||
static void update_attribute_element_size(Mesh *mesh,
|
||||
@ -1554,9 +1554,9 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
|
||||
}
|
||||
}
|
||||
|
||||
dscene->attributes_float.resize(attr_float_size);
|
||||
dscene->attributes_float3.resize(attr_float3_size);
|
||||
dscene->attributes_uchar4.resize(attr_uchar4_size);
|
||||
dscene->attributes_float.alloc(attr_float_size);
|
||||
dscene->attributes_float3.alloc(attr_float3_size);
|
||||
dscene->attributes_uchar4.alloc(attr_uchar4_size);
|
||||
|
||||
size_t attr_float_offset = 0;
|
||||
size_t attr_float3_offset = 0;
|
||||
@ -1617,13 +1617,13 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
|
||||
progress.set_status("Updating Mesh", "Copying Attributes to device");
|
||||
|
||||
if(dscene->attributes_float.size()) {
|
||||
device->tex_alloc(dscene->attributes_float);
|
||||
dscene->attributes_float.copy_to_device();
|
||||
}
|
||||
if(dscene->attributes_float3.size()) {
|
||||
device->tex_alloc(dscene->attributes_float3);
|
||||
dscene->attributes_float3.copy_to_device();
|
||||
}
|
||||
if(dscene->attributes_uchar4.size()) {
|
||||
device->tex_alloc(dscene->attributes_uchar4);
|
||||
dscene->attributes_uchar4.copy_to_device();
|
||||
}
|
||||
}
|
||||
|
||||
@ -1671,7 +1671,7 @@ void MeshManager::mesh_calc_offset(Scene *scene)
|
||||
}
|
||||
}
|
||||
|
||||
void MeshManager::device_update_mesh(Device *device,
|
||||
void MeshManager::device_update_mesh(Device *,
|
||||
DeviceScene *dscene,
|
||||
Scene *scene,
|
||||
bool for_displacement,
|
||||
@ -1732,11 +1732,11 @@ void MeshManager::device_update_mesh(Device *device,
|
||||
/* normals */
|
||||
progress.set_status("Updating Mesh", "Computing normals");
|
||||
|
||||
uint *tri_shader = dscene->tri_shader.resize(tri_size);
|
||||
float4 *vnormal = dscene->tri_vnormal.resize(vert_size);
|
||||
uint4 *tri_vindex = dscene->tri_vindex.resize(tri_size);
|
||||
uint *tri_patch = dscene->tri_patch.resize(tri_size);
|
||||
float2 *tri_patch_uv = dscene->tri_patch_uv.resize(vert_size);
|
||||
uint *tri_shader = dscene->tri_shader.alloc(tri_size);
|
||||
float4 *vnormal = dscene->tri_vnormal.alloc(vert_size);
|
||||
uint4 *tri_vindex = dscene->tri_vindex.alloc(tri_size);
|
||||
uint *tri_patch = dscene->tri_patch.alloc(tri_size);
|
||||
float2 *tri_patch_uv = dscene->tri_patch_uv.alloc(vert_size);
|
||||
|
||||
foreach(Mesh *mesh, scene->meshes) {
|
||||
mesh->pack_normals(scene,
|
||||
@ -1754,32 +1754,32 @@ void MeshManager::device_update_mesh(Device *device,
|
||||
/* vertex coordinates */
|
||||
progress.set_status("Updating Mesh", "Copying Mesh to device");
|
||||
|
||||
device->tex_alloc(dscene->tri_shader);
|
||||
device->tex_alloc(dscene->tri_vnormal);
|
||||
device->tex_alloc(dscene->tri_vindex);
|
||||
device->tex_alloc(dscene->tri_patch);
|
||||
device->tex_alloc(dscene->tri_patch_uv);
|
||||
dscene->tri_shader.copy_to_device();
|
||||
dscene->tri_vnormal.copy_to_device();
|
||||
dscene->tri_vindex.copy_to_device();
|
||||
dscene->tri_patch.copy_to_device();
|
||||
dscene->tri_patch_uv.copy_to_device();
|
||||
}
|
||||
|
||||
if(curve_size != 0) {
|
||||
progress.set_status("Updating Mesh", "Copying Strands to device");
|
||||
|
||||
float4 *curve_keys = dscene->curve_keys.resize(curve_key_size);
|
||||
float4 *curves = dscene->curves.resize(curve_size);
|
||||
float4 *curve_keys = dscene->curve_keys.alloc(curve_key_size);
|
||||
float4 *curves = dscene->curves.alloc(curve_size);
|
||||
|
||||
foreach(Mesh *mesh, scene->meshes) {
|
||||
mesh->pack_curves(scene, &curve_keys[mesh->curvekey_offset], &curves[mesh->curve_offset], mesh->curvekey_offset);
|
||||
if(progress.get_cancel()) return;
|
||||
}
|
||||
|
||||
device->tex_alloc(dscene->curve_keys);
|
||||
device->tex_alloc(dscene->curves);
|
||||
dscene->curve_keys.copy_to_device();
|
||||
dscene->curves.copy_to_device();
|
||||
}
|
||||
|
||||
if(patch_size != 0) {
|
||||
progress.set_status("Updating Mesh", "Copying Patches to device");
|
||||
|
||||
uint *patch_data = dscene->patches.resize(patch_size);
|
||||
uint *patch_data = dscene->patches.alloc(patch_size);
|
||||
|
||||
foreach(Mesh *mesh, scene->meshes) {
|
||||
mesh->pack_patches(&patch_data[mesh->patch_offset], mesh->vert_offset, mesh->face_offset, mesh->corner_offset);
|
||||
@ -1791,11 +1791,11 @@ void MeshManager::device_update_mesh(Device *device,
|
||||
if(progress.get_cancel()) return;
|
||||
}
|
||||
|
||||
device->tex_alloc(dscene->patches);
|
||||
dscene->patches.copy_to_device();
|
||||
}
|
||||
|
||||
if(for_displacement) {
|
||||
float4 *prim_tri_verts = dscene->prim_tri_verts.resize(tri_size * 3);
|
||||
float4 *prim_tri_verts = dscene->prim_tri_verts.alloc(tri_size * 3);
|
||||
foreach(Mesh *mesh, scene->meshes) {
|
||||
for(size_t i = 0; i < mesh->num_triangles(); ++i) {
|
||||
Mesh::Triangle t = mesh->get_triangle(i);
|
||||
@ -1805,7 +1805,7 @@ void MeshManager::device_update_mesh(Device *device,
|
||||
prim_tri_verts[offset + 2] = float3_to_float4(mesh->verts[t.v[2]]);
|
||||
}
|
||||
}
|
||||
device->tex_alloc(dscene->prim_tri_verts);
|
||||
dscene->prim_tri_verts.copy_to_device();
|
||||
}
|
||||
}
|
||||
|
||||
@ -1841,43 +1841,43 @@ void MeshManager::device_update_bvh(Device *device, DeviceScene *dscene, Scene *
|
||||
|
||||
if(pack.nodes.size()) {
|
||||
dscene->bvh_nodes.steal_data(pack.nodes);
|
||||
device->tex_alloc(dscene->bvh_nodes);
|
||||
dscene->bvh_nodes.copy_to_device();
|
||||
}
|
||||
if(pack.leaf_nodes.size()) {
|
||||
dscene->bvh_leaf_nodes.steal_data(pack.leaf_nodes);
|
||||
device->tex_alloc(dscene->bvh_leaf_nodes);
|
||||
dscene->bvh_leaf_nodes.copy_to_device();
|
||||
}
|
||||
if(pack.object_node.size()) {
|
||||
dscene->object_node.steal_data(pack.object_node);
|
||||
device->tex_alloc(dscene->object_node);
|
||||
dscene->object_node.copy_to_device();
|
||||
}
|
||||
if(pack.prim_tri_index.size()) {
|
||||
dscene->prim_tri_index.steal_data(pack.prim_tri_index);
|
||||
device->tex_alloc(dscene->prim_tri_index);
|
||||
dscene->prim_tri_index.copy_to_device();
|
||||
}
|
||||
if(pack.prim_tri_verts.size()) {
|
||||
dscene->prim_tri_verts.steal_data(pack.prim_tri_verts);
|
||||
device->tex_alloc(dscene->prim_tri_verts);
|
||||
dscene->prim_tri_verts.copy_to_device();
|
||||
}
|
||||
if(pack.prim_type.size()) {
|
||||
dscene->prim_type.steal_data(pack.prim_type);
|
||||
device->tex_alloc(dscene->prim_type);
|
||||
dscene->prim_type.copy_to_device();
|
||||
}
|
||||
if(pack.prim_visibility.size()) {
|
||||
dscene->prim_visibility.steal_data(pack.prim_visibility);
|
||||
device->tex_alloc(dscene->prim_visibility);
|
||||
dscene->prim_visibility.copy_to_device();
|
||||
}
|
||||
if(pack.prim_index.size()) {
|
||||
dscene->prim_index.steal_data(pack.prim_index);
|
||||
device->tex_alloc(dscene->prim_index);
|
||||
dscene->prim_index.copy_to_device();
|
||||
}
|
||||
if(pack.prim_object.size()) {
|
||||
dscene->prim_object.steal_data(pack.prim_object);
|
||||
device->tex_alloc(dscene->prim_object);
|
||||
dscene->prim_object.copy_to_device();
|
||||
}
|
||||
if(pack.prim_time.size()) {
|
||||
dscene->prim_time.steal_data(pack.prim_time);
|
||||
device->tex_alloc(dscene->prim_time);
|
||||
dscene->prim_time.copy_to_device();
|
||||
}
|
||||
|
||||
dscene->data.bvh.root = pack.root_index;
|
||||
@ -2142,51 +2142,28 @@ void MeshManager::device_update(Device *device, DeviceScene *dscene, Scene *scen
|
||||
|
||||
void MeshManager::device_free(Device *device, DeviceScene *dscene)
|
||||
{
|
||||
device->tex_free(dscene->bvh_nodes);
|
||||
device->tex_free(dscene->bvh_leaf_nodes);
|
||||
device->tex_free(dscene->object_node);
|
||||
device->tex_free(dscene->prim_tri_verts);
|
||||
device->tex_free(dscene->prim_tri_index);
|
||||
device->tex_free(dscene->prim_type);
|
||||
device->tex_free(dscene->prim_visibility);
|
||||
device->tex_free(dscene->prim_index);
|
||||
device->tex_free(dscene->prim_object);
|
||||
device->tex_free(dscene->prim_time);
|
||||
device->tex_free(dscene->tri_shader);
|
||||
device->tex_free(dscene->tri_vnormal);
|
||||
device->tex_free(dscene->tri_vindex);
|
||||
device->tex_free(dscene->tri_patch);
|
||||
device->tex_free(dscene->tri_patch_uv);
|
||||
device->tex_free(dscene->curves);
|
||||
device->tex_free(dscene->curve_keys);
|
||||
device->tex_free(dscene->patches);
|
||||
device->tex_free(dscene->attributes_map);
|
||||
device->tex_free(dscene->attributes_float);
|
||||
device->tex_free(dscene->attributes_float3);
|
||||
device->tex_free(dscene->attributes_uchar4);
|
||||
|
||||
dscene->bvh_nodes.clear();
|
||||
dscene->bvh_leaf_nodes.clear();
|
||||
dscene->object_node.clear();
|
||||
dscene->prim_tri_verts.clear();
|
||||
dscene->prim_tri_index.clear();
|
||||
dscene->prim_type.clear();
|
||||
dscene->prim_visibility.clear();
|
||||
dscene->prim_index.clear();
|
||||
dscene->prim_object.clear();
|
||||
dscene->prim_time.clear();
|
||||
dscene->tri_shader.clear();
|
||||
dscene->tri_vnormal.clear();
|
||||
dscene->tri_vindex.clear();
|
||||
dscene->tri_patch.clear();
|
||||
dscene->tri_patch_uv.clear();
|
||||
dscene->curves.clear();
|
||||
dscene->curve_keys.clear();
|
||||
dscene->patches.clear();
|
||||
dscene->attributes_map.clear();
|
||||
dscene->attributes_float.clear();
|
||||
dscene->attributes_float3.clear();
|
||||
dscene->attributes_uchar4.clear();
|
||||
dscene->bvh_nodes.free();
|
||||
dscene->bvh_leaf_nodes.free();
|
||||
dscene->object_node.free();
|
||||
dscene->prim_tri_verts.free();
|
||||
dscene->prim_tri_index.free();
|
||||
dscene->prim_type.free();
|
||||
dscene->prim_visibility.free();
|
||||
dscene->prim_index.free();
|
||||
dscene->prim_object.free();
|
||||
dscene->prim_time.free();
|
||||
dscene->tri_shader.free();
|
||||
dscene->tri_vnormal.free();
|
||||
dscene->tri_vindex.free();
|
||||
dscene->tri_patch.free();
|
||||
dscene->tri_patch_uv.free();
|
||||
dscene->curves.free();
|
||||
dscene->curve_keys.free();
|
||||
dscene->patches.free();
|
||||
dscene->attributes_map.free();
|
||||
dscene->attributes_float.free();
|
||||
dscene->attributes_float3.free();
|
||||
dscene->attributes_uchar4.free();
|
||||
|
||||
#ifdef WITH_OSL
|
||||
OSLGlobals *og = (OSLGlobals*)device->osl_memory();
|
||||
|
@ -65,7 +65,7 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
|
||||
const size_t num_verts = mesh->verts.size();
|
||||
vector<bool> done(num_verts, false);
|
||||
device_vector<uint4> d_input(device, "displace_input", MEM_READ_ONLY);
|
||||
uint4 *d_input_data = d_input.resize(num_verts);
|
||||
uint4 *d_input_data = d_input.alloc(num_verts);
|
||||
size_t d_input_size = 0;
|
||||
|
||||
size_t num_triangles = mesh->num_triangles();
|
||||
@ -116,16 +116,13 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
|
||||
|
||||
/* run device task */
|
||||
device_vector<float4> d_output(device, "displace_output", MEM_WRITE_ONLY);
|
||||
d_output.resize(d_input_size);
|
||||
d_output.alloc(d_input_size);
|
||||
d_output.zero_to_device();
|
||||
d_input.copy_to_device();
|
||||
|
||||
/* needs to be up to data for attribute access */
|
||||
device->const_copy_to("__data", &dscene->data, sizeof(dscene->data));
|
||||
|
||||
device->mem_alloc(d_input);
|
||||
device->mem_copy_to(d_input);
|
||||
device->mem_alloc(d_output);
|
||||
device->mem_zero(d_output);
|
||||
|
||||
DeviceTask task(DeviceTask::SHADER);
|
||||
task.shader_input = d_input.device_pointer;
|
||||
task.shader_output = d_output.device_pointer;
|
||||
@ -139,14 +136,13 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
|
||||
device->task_wait();
|
||||
|
||||
if(progress.get_cancel()) {
|
||||
device->mem_free(d_input);
|
||||
device->mem_free(d_output);
|
||||
d_input.free();
|
||||
d_output.free();
|
||||
return false;
|
||||
}
|
||||
|
||||
device->mem_copy_from(d_output, 0, 1, d_output.size(), sizeof(float4));
|
||||
device->mem_free(d_input);
|
||||
device->mem_free(d_output);
|
||||
d_output.copy_from_device(0, 1, d_output.size());
|
||||
d_input.free();
|
||||
|
||||
/* read result */
|
||||
done.clear();
|
||||
@ -183,6 +179,8 @@ bool MeshManager::displace(Device *device, DeviceScene *dscene, Scene *scene, Me
|
||||
}
|
||||
}
|
||||
|
||||
d_output.free();
|
||||
|
||||
/* for displacement method both, we only need to recompute the face
|
||||
* normals, as bump mapping in the shader will already alter the
|
||||
* vertex normal, so we start from the non-displaced vertex normals
|
||||
|
@ -488,9 +488,9 @@ void ObjectManager::device_update_transforms(Device *device,
|
||||
state.queue_start_object = 0;
|
||||
|
||||
state.object_flag = object_flag;
|
||||
state.objects = dscene->objects.resize(OBJECT_SIZE*scene->objects.size());
|
||||
state.objects = dscene->objects.alloc(OBJECT_SIZE*scene->objects.size());
|
||||
if(state.need_motion == Scene::MOTION_PASS) {
|
||||
state.objects_vector = dscene->objects_vector.resize(OBJECT_VECTOR_SIZE*scene->objects.size());
|
||||
state.objects_vector = dscene->objects_vector.alloc(OBJECT_VECTOR_SIZE*scene->objects.size());
|
||||
}
|
||||
else {
|
||||
state.objects_vector = NULL;
|
||||
@ -534,9 +534,9 @@ void ObjectManager::device_update_transforms(Device *device,
|
||||
}
|
||||
}
|
||||
|
||||
device->tex_alloc(dscene->objects);
|
||||
dscene->objects.copy_to_device();
|
||||
if(state.need_motion == Scene::MOTION_PASS) {
|
||||
device->tex_alloc(dscene->objects_vector);
|
||||
dscene->objects_vector.copy_to_device();
|
||||
}
|
||||
|
||||
dscene->data.bvh.have_motion = state.have_motion;
|
||||
@ -557,7 +557,7 @@ void ObjectManager::device_update(Device *device, DeviceScene *dscene, Scene *sc
|
||||
return;
|
||||
|
||||
/* object info flag */
|
||||
uint *object_flag = dscene->object_flag.resize(scene->objects.size());
|
||||
uint *object_flag = dscene->object_flag.alloc(scene->objects.size());
|
||||
|
||||
/* set object transform matrices, before applying static transforms */
|
||||
progress.set_status("Updating Objects", "Copying Transformations to device");
|
||||
@ -573,7 +573,7 @@ void ObjectManager::device_update(Device *device, DeviceScene *dscene, Scene *sc
|
||||
}
|
||||
}
|
||||
|
||||
void ObjectManager::device_update_flags(Device *device,
|
||||
void ObjectManager::device_update_flags(Device *,
|
||||
DeviceScene *dscene,
|
||||
Scene *scene,
|
||||
Progress& /*progress*/,
|
||||
@ -638,10 +638,10 @@ void ObjectManager::device_update_flags(Device *device,
|
||||
}
|
||||
|
||||
/* allocate object flag */
|
||||
device->tex_alloc(dscene->object_flag);
|
||||
dscene->object_flag.copy_to_device();
|
||||
}
|
||||
|
||||
void ObjectManager::device_update_patch_map_offsets(Device *device, DeviceScene *dscene, Scene *scene)
|
||||
void ObjectManager::device_update_patch_map_offsets(Device *, DeviceScene *dscene, Scene *scene)
|
||||
{
|
||||
if(scene->objects.size() == 0) {
|
||||
return;
|
||||
@ -671,21 +671,15 @@ void ObjectManager::device_update_patch_map_offsets(Device *device, DeviceScene
|
||||
}
|
||||
|
||||
if(update) {
|
||||
device->tex_free(dscene->objects);
|
||||
device->tex_alloc(dscene->objects);
|
||||
dscene->objects.copy_to_device();
|
||||
}
|
||||
}
|
||||
|
||||
void ObjectManager::device_free(Device *device, DeviceScene *dscene)
|
||||
void ObjectManager::device_free(Device *, DeviceScene *dscene)
|
||||
{
|
||||
device->tex_free(dscene->objects);
|
||||
dscene->objects.clear();
|
||||
|
||||
device->tex_free(dscene->objects_vector);
|
||||
dscene->objects_vector.clear();
|
||||
|
||||
device->tex_free(dscene->object_flag);
|
||||
dscene->object_flag.clear();
|
||||
dscene->objects.free();
|
||||
dscene->objects_vector.free();
|
||||
dscene->object_flag.free();
|
||||
}
|
||||
|
||||
void ObjectManager::apply_static_transforms(DeviceScene *dscene, Scene *scene, uint *object_flag, Progress& progress)
|
||||
|
@ -52,7 +52,7 @@ ParticleSystemManager::~ParticleSystemManager()
|
||||
{
|
||||
}
|
||||
|
||||
void ParticleSystemManager::device_update_particles(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
|
||||
void ParticleSystemManager::device_update_particles(Device *, DeviceScene *dscene, Scene *scene, Progress& progress)
|
||||
{
|
||||
/* count particles.
|
||||
* adds one dummy particle at the beginning to avoid invalid lookups,
|
||||
@ -61,7 +61,7 @@ void ParticleSystemManager::device_update_particles(Device *device, DeviceScene
|
||||
for(size_t j = 0; j < scene->particle_systems.size(); j++)
|
||||
num_particles += scene->particle_systems[j]->particles.size();
|
||||
|
||||
float4 *particles = dscene->particles.resize(PARTICLE_SIZE*num_particles);
|
||||
float4 *particles = dscene->particles.alloc(PARTICLE_SIZE*num_particles);
|
||||
|
||||
/* dummy particle */
|
||||
particles[0] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
|
||||
@ -91,7 +91,7 @@ void ParticleSystemManager::device_update_particles(Device *device, DeviceScene
|
||||
}
|
||||
}
|
||||
|
||||
device->tex_alloc(dscene->particles);
|
||||
dscene->particles.copy_to_device();
|
||||
}
|
||||
|
||||
void ParticleSystemManager::device_update(Device *device, DeviceScene *dscene, Scene *scene, Progress& progress)
|
||||
@ -112,10 +112,9 @@ void ParticleSystemManager::device_update(Device *device, DeviceScene *dscene, S
|
||||
need_update = false;
|
||||
}
|
||||
|
||||
void ParticleSystemManager::device_free(Device *device, DeviceScene *dscene)
|
||||
void ParticleSystemManager::device_free(Device *, DeviceScene *dscene)
|
||||
{
|
||||
device->tex_free(dscene->particles);
|
||||
dscene->particles.clear();
|
||||
dscene->particles.free();
|
||||
}
|
||||
|
||||
void ParticleSystemManager::tag_update(Scene * /*scene*/)
|
||||
|
@ -41,40 +41,40 @@
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
DeviceScene::DeviceScene(Device *device)
|
||||
: bvh_nodes(device, "__bvh_nodes"),
|
||||
bvh_leaf_nodes(device, "__bvh_leaf_nodes"),
|
||||
object_node(device, "__object_node"),
|
||||
prim_tri_index(device, "__prim_tri_index"),
|
||||
prim_tri_verts(device, "__prim_tri_verts"),
|
||||
prim_type(device, "__prim_type"),
|
||||
prim_visibility(device, "__prim_visibility"),
|
||||
prim_index(device, "__prim_index"),
|
||||
prim_object(device, "__prim_object"),
|
||||
prim_time(device, "__prim_time"),
|
||||
tri_shader(device, "__tri_shader"),
|
||||
tri_vnormal(device, "__tri_vnormal"),
|
||||
tri_vindex(device, "__tri_vindex"),
|
||||
tri_patch(device, "__tri_patch"),
|
||||
tri_patch_uv(device, "__tri_patch_uv"),
|
||||
curves(device, "__curves"),
|
||||
curve_keys(device, "__curve_keys"),
|
||||
patches(device, "__patches"),
|
||||
objects(device, "__objects"),
|
||||
objects_vector(device, "__objects_vector"),
|
||||
attributes_map(device, "__attributes_map"),
|
||||
attributes_float(device, "__attributes_float"),
|
||||
attributes_float3(device, "__attributes_float3"),
|
||||
attributes_uchar4(device, "__attributes_uchar4"),
|
||||
light_distribution(device, "__light_distribution"),
|
||||
light_data(device, "__light_data"),
|
||||
light_background_marginal_cdf(device, "__light_background_marginal_cdf"),
|
||||
light_background_conditional_cdf(device, "__light_background_conditional_cdf"),
|
||||
particles(device, "__particles"),
|
||||
svm_nodes(device, "__svm_nodes"),
|
||||
shader_flag(device, "__shader_flag"),
|
||||
object_flag(device, "__object_flag"),
|
||||
lookup_table(device, "__lookup_table"),
|
||||
sobol_directions(device, "__sobol_directions")
|
||||
: bvh_nodes(device, "__bvh_nodes", MEM_TEXTURE),
|
||||
bvh_leaf_nodes(device, "__bvh_leaf_nodes", MEM_TEXTURE),
|
||||
object_node(device, "__object_node", MEM_TEXTURE),
|
||||
prim_tri_index(device, "__prim_tri_index", MEM_TEXTURE),
|
||||
prim_tri_verts(device, "__prim_tri_verts", MEM_TEXTURE),
|
||||
prim_type(device, "__prim_type", MEM_TEXTURE),
|
||||
prim_visibility(device, "__prim_visibility", MEM_TEXTURE),
|
||||
prim_index(device, "__prim_index", MEM_TEXTURE),
|
||||
prim_object(device, "__prim_object", MEM_TEXTURE),
|
||||
prim_time(device, "__prim_time", MEM_TEXTURE),
|
||||
tri_shader(device, "__tri_shader", MEM_TEXTURE),
|
||||
tri_vnormal(device, "__tri_vnormal", MEM_TEXTURE),
|
||||
tri_vindex(device, "__tri_vindex", MEM_TEXTURE),
|
||||
tri_patch(device, "__tri_patch", MEM_TEXTURE),
|
||||
tri_patch_uv(device, "__tri_patch_uv", MEM_TEXTURE),
|
||||
curves(device, "__curves", MEM_TEXTURE),
|
||||
curve_keys(device, "__curve_keys", MEM_TEXTURE),
|
||||
patches(device, "__patches", MEM_TEXTURE),
|
||||
objects(device, "__objects", MEM_TEXTURE),
|
||||
objects_vector(device, "__objects_vector", MEM_TEXTURE),
|
||||
attributes_map(device, "__attributes_map", MEM_TEXTURE),
|
||||
attributes_float(device, "__attributes_float", MEM_TEXTURE),
|
||||
attributes_float3(device, "__attributes_float3", MEM_TEXTURE),
|
||||
attributes_uchar4(device, "__attributes_uchar4", MEM_TEXTURE),
|
||||
light_distribution(device, "__light_distribution", MEM_TEXTURE),
|
||||
light_data(device, "__light_data", MEM_TEXTURE),
|
||||
light_background_marginal_cdf(device, "__light_background_marginal_cdf", MEM_TEXTURE),
|
||||
light_background_conditional_cdf(device, "__light_background_conditional_cdf", MEM_TEXTURE),
|
||||
particles(device, "__particles", MEM_TEXTURE),
|
||||
svm_nodes(device, "__svm_nodes", MEM_TEXTURE),
|
||||
shader_flag(device, "__shader_flag", MEM_TEXTURE),
|
||||
object_flag(device, "__object_flag", MEM_TEXTURE),
|
||||
lookup_table(device, "__lookup_table", MEM_TEXTURE),
|
||||
sobol_directions(device, "__sobol_directions", MEM_TEXTURE)
|
||||
{
|
||||
memset(&data, 0, sizeof(data));
|
||||
}
|
||||
|
@ -106,11 +106,11 @@ Session::~Session()
|
||||
delete display;
|
||||
|
||||
display = new DisplayBuffer(device, false);
|
||||
display->reset(device, buffers->params);
|
||||
display->reset(buffers->params);
|
||||
tonemap(params.samples);
|
||||
|
||||
progress.set_status("Writing Image", params.output_path);
|
||||
display->write(device, params.output_path);
|
||||
display->write(params.output_path);
|
||||
}
|
||||
|
||||
/* clean up */
|
||||
@ -399,7 +399,7 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile)
|
||||
|
||||
/* allocate buffers */
|
||||
tile->buffers = new RenderBuffers(tile_device);
|
||||
tile->buffers->reset(tile_device, buffer_params);
|
||||
tile->buffers->reset(buffer_params);
|
||||
}
|
||||
|
||||
tile->buffers->params.get_offset_stride(rtile.offset, rtile.stride);
|
||||
@ -756,9 +756,9 @@ void Session::reset_(BufferParams& buffer_params, int samples)
|
||||
{
|
||||
if(buffers && buffer_params.modified(tile_manager.params)) {
|
||||
gpu_draw_ready = false;
|
||||
buffers->reset(device, buffer_params);
|
||||
buffers->reset(buffer_params);
|
||||
if(display) {
|
||||
display->reset(device, buffer_params);
|
||||
display->reset(buffer_params);
|
||||
}
|
||||
}
|
||||
|
||||
@ -923,7 +923,7 @@ void Session::render()
|
||||
{
|
||||
/* Clear buffers. */
|
||||
if(buffers && tile_manager.state.sample == tile_manager.range_start_sample) {
|
||||
buffers->zero(device);
|
||||
buffers->zero();
|
||||
}
|
||||
|
||||
/* Add path trace task. */
|
||||
|
@ -416,14 +416,13 @@ void ShaderManager::device_update_common(Device *device,
|
||||
Scene *scene,
|
||||
Progress& /*progress*/)
|
||||
{
|
||||
device->tex_free(dscene->shader_flag);
|
||||
dscene->shader_flag.clear();
|
||||
dscene->shader_flag.free();
|
||||
|
||||
if(scene->shaders.size() == 0)
|
||||
return;
|
||||
|
||||
uint shader_flag_size = scene->shaders.size()*SHADER_SIZE;
|
||||
uint *shader_flag = dscene->shader_flag.resize(shader_flag_size);
|
||||
uint *shader_flag = dscene->shader_flag.alloc(shader_flag_size);
|
||||
uint i = 0;
|
||||
bool has_volumes = false;
|
||||
bool has_transparent_shadow = false;
|
||||
@ -479,7 +478,7 @@ void ShaderManager::device_update_common(Device *device,
|
||||
has_transparent_shadow |= (flag & SD_HAS_TRANSPARENT_SHADOW) != 0;
|
||||
}
|
||||
|
||||
device->tex_alloc(dscene->shader_flag);
|
||||
dscene->shader_flag.copy_to_device();
|
||||
|
||||
/* lookup tables */
|
||||
KernelTables *ktables = &dscene->data.tables;
|
||||
@ -504,12 +503,11 @@ void ShaderManager::device_update_common(Device *device,
|
||||
kintegrator->transparent_shadows = has_transparent_shadow;
|
||||
}
|
||||
|
||||
void ShaderManager::device_free_common(Device *device, DeviceScene *dscene, Scene *scene)
|
||||
void ShaderManager::device_free_common(Device *, DeviceScene *dscene, Scene *scene)
|
||||
{
|
||||
scene->lookup_tables->remove_table(&beckmann_table_offset);
|
||||
|
||||
device->tex_free(dscene->shader_flag);
|
||||
dscene->shader_flag.clear();
|
||||
dscene->shader_flag.free();
|
||||
}
|
||||
|
||||
void ShaderManager::add_default(Scene *scene)
|
||||
|
@ -130,7 +130,7 @@ void SVMShaderManager::device_update(Device *device, DeviceScene *dscene, Scene
|
||||
}
|
||||
|
||||
dscene->svm_nodes.steal_data(svm_nodes);
|
||||
device->tex_alloc(dscene->svm_nodes);
|
||||
dscene->svm_nodes.copy_to_device();
|
||||
|
||||
for(i = 0; i < scene->shaders.size(); i++) {
|
||||
Shader *shader = scene->shaders[i];
|
||||
@ -150,8 +150,7 @@ void SVMShaderManager::device_free(Device *device, DeviceScene *dscene, Scene *s
|
||||
{
|
||||
device_free_common(device, dscene, scene);
|
||||
|
||||
device->tex_free(dscene->svm_nodes);
|
||||
dscene->svm_nodes.clear();
|
||||
dscene->svm_nodes.free();
|
||||
}
|
||||
|
||||
/* Graph Compiler */
|
||||
|
@ -35,25 +35,22 @@ LookupTables::~LookupTables()
|
||||
assert(lookup_tables.size() == 0);
|
||||
}
|
||||
|
||||
void LookupTables::device_update(Device *device, DeviceScene *dscene)
|
||||
void LookupTables::device_update(Device *, DeviceScene *dscene)
|
||||
{
|
||||
if(!need_update)
|
||||
return;
|
||||
|
||||
VLOG(1) << "Total " << lookup_tables.size() << " lookup tables.";
|
||||
|
||||
device->tex_free(dscene->lookup_table);
|
||||
|
||||
if(lookup_tables.size() > 0)
|
||||
device->tex_alloc(dscene->lookup_table);
|
||||
dscene->lookup_table.copy_to_device();
|
||||
|
||||
need_update = false;
|
||||
}
|
||||
|
||||
void LookupTables::device_free(Device *device, DeviceScene *dscene)
|
||||
void LookupTables::device_free(Device *, DeviceScene *dscene)
|
||||
{
|
||||
device->tex_free(dscene->lookup_table);
|
||||
dscene->lookup_table.clear();
|
||||
dscene->lookup_table.free();
|
||||
}
|
||||
|
||||
static size_t round_up_to_multiple(size_t size, size_t chunk)
|
||||
|
@ -177,6 +177,14 @@ public:
|
||||
}
|
||||
}
|
||||
|
||||
T *steal_pointer()
|
||||
{
|
||||
T *ptr = data_;
|
||||
data_ = NULL;
|
||||
clear();
|
||||
return ptr;
|
||||
}
|
||||
|
||||
T* resize(size_t newsize)
|
||||
{
|
||||
if(newsize == 0) {
|
||||
|
Loading…
Reference in New Issue
Block a user