diff options
Diffstat (limited to 'intern')
30 files changed, 934 insertions, 855 deletions
diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp index 9de10c184fb..41fbe7ce81b 100644 --- a/intern/cycles/device/device.cpp +++ b/intern/cycles/device/device.cpp @@ -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); diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h index 6bb65cde2a3..316bf70a5c3 100644 --- a/intern/cycles/device/device.h +++ b/intern/cycles/device/device.h @@ -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; diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp index b4398f21014..32ab18fe164 100644 --- a/intern/cycles/device/device_cpu.cpp +++ b/intern/cycles/device/device_cpu.cpp @@ -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; } diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp index be606a92434..aa6386e455b 100644 --- a/intern/cycles/device/device_cuda.cpp +++ b/intern/cycles/device/device_cuda.cpp @@ -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; - - pmem.w = mem.data_width; - pmem.h = mem.data_height; + PixelMem pmem; - CUDAContextScope scope(this); + pmem.w = mem.data_width; + pmem.h = mem.data_height; - 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); + CUDAContextScope scope(this); - glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + 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); - 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); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); - CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE); + 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); - if(result == CUDA_SUCCESS) { - mem.device_pointer = pmem.cuTexId; - pixel_mem_map[mem.device_pointer] = pmem; + CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE); - mem.device_size = mem.memory_size(); - stats.mem_alloc(mem.device_size); + if(result == CUDA_SUCCESS) { + mem.device_pointer = pmem.cuTexId; + pixel_mem_map[mem.device_pointer] = pmem; - return; - } - else { - /* failed to register buffer, fallback to no interop */ - glDeleteBuffers(1, &pmem.cuPBO); - glDeleteTextures(1, &pmem.cuTexId); + mem.device_size = mem.memory_size(); + stats.mem_alloc(mem.device_size); - 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]; - - 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); + PixelMem pmem = pixel_mem_map[mem.device_pointer]; - return; - } + CUDAContextScope scope(this); - 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]; - - CUDAContextScope scope(this); - - cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource)); - glDeleteBuffers(1, &pmem.cuPBO); - glDeleteTextures(1, &pmem.cuTexId); + PixelMem pmem = pixel_mem_map[mem.device_pointer]; - pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer)); - mem.device_pointer = 0; + CUDAContextScope scope(this); - stats.mem_free(mem.device_size); - mem.device_size = 0; + cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource)); + glDeleteBuffers(1, &pmem.cuPBO); + glDeleteTextures(1, &pmem.cuTexId); - return; - } + pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer)); + mem.device_pointer = 0; - 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; } diff --git a/intern/cycles/device/device_denoising.cpp b/intern/cycles/device/device_denoising.cpp index 2c3bfefd8b0..2d39721e3d3 100644 --- a/intern/cycles/device/device_denoising.cpp +++ b/intern/cycles/device/device_denoising.cpp @@ -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; } diff --git a/intern/cycles/device/device_memory.cpp b/intern/cycles/device/device_memory.cpp index 98fa638ef8e..9f4f60e7531 100644 --- a/intern/cycles/device/device_memory.cpp +++ b/intern/cycles/device/device_memory.cpp @@ -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) diff --git a/intern/cycles/device/device_memory.h b/intern/cycles/device/device_memory.h index 3dfecde59d8..7bf8bdc1cea 100644 --- a/intern/cycles/device/device_memory.h +++ b/intern/cycles/device/device_memory.h @@ -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() + { + free(); + } + + void alloc_to_device(size_t num) + { + data_size = num*sizeof(T); + device_alloc(); + } + + void free() + { + device_free(); + } + + void zero_to_device() { - device_memory::resize(num*sizeof(T)); + 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() {} + virtual ~device_vector() + { + free(); + } - /* vector functions */ - T *resize(size_t width, size_t height = 0, size_t depth = 0) + /* Host memory allocation. */ + T *alloc(size_t width, size_t height = 0, size_t depth = 0) { - data_size = width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth); - if(data.resize(data_size) == NULL) { - clear(); - return NULL; + 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; + 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_pointer = (device_ptr)&data[0]; - return &data[0]; + + 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; + device_free(); + host_free(data_pointer, sizeof(T)*data_size); + + data_size = from.size(); + data_width = 0; data_height = 0; data_depth = 0; + data_pointer = (device_ptr)from.steal_pointer(); + assert(device_pointer == 0); } - void clear() + /* Free device and host memory. */ + void free() { - data.clear(); - data_pointer = 0; + device_free(); + host_free(data_pointer, sizeof(T)*data_size); + + data_size = 0; data_width = 0; data_height = 0; data_depth = 0; - data_size = 0; - device_pointer = 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; }; diff --git a/intern/cycles/device/device_multi.cpp b/intern/cycles/device/device_multi.cpp index 7f7fbc0d1d3..0a6dd90c86d 100644 --- a/intern/cycles/device/device_multi.cpp +++ b/intern/cycles/device/device_multi.cpp @@ -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(); } } } diff --git a/intern/cycles/device/device_network.cpp b/intern/cycles/device/device_network.cpp index bdc88b6acae..fa231c817e6 100644 --- a/intern/cycles/device/device_network.cpp +++ b/intern/cycles/device/device_network.cpp @@ -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); - - size_t data_size = mem.memory_size(); + 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]; - /* get pointer to memory buffer for device buffer */ - 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; + } - /* 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); diff --git a/intern/cycles/device/device_network.h b/intern/cycles/device/device_network.h index 8a53290f421..a38d962c0af 100644 --- a/intern/cycles/device/device_network.h +++ b/intern/cycles/device/device_network.h @@ -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) diff --git a/intern/cycles/device/device_split_kernel.cpp b/intern/cycles/device/device_split_kernel.cpp index 6c8befa89be..f2839a8b1b9 100644 --- a/intern/cycles/device/device_split_kernel.cpp +++ b/intern/cycles/device/device_split_kernel.cpp @@ -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; diff --git a/intern/cycles/device/opencl/memory_manager.cpp b/intern/cycles/device/opencl/memory_manager.cpp index e48367b8987..a791b374774 100644 --- a/intern/cycles/device/opencl/memory_manager.cpp +++ b/intern/cycles/device/opencl/memory_manager.cpp @@ -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() diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp index 90f461b4c98..5e9debc3b17 100644 --- a/intern/cycles/device/opencl/opencl_base.cpp +++ b/intern/cycles/device/opencl/opencl_base.cpp @@ -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); diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp index c966ebe0c5e..96139afa450 100644 --- a/intern/cycles/device/opencl/opencl_split.cpp +++ b/intern/cycles/device/opencl/opencl_split.cpp @@ -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()", diff --git a/intern/cycles/render/bake.cpp b/intern/cycles/render/bake.cpp index 66615bf336c..99f68b6aa00 100644 --- a/intern/cycles/render/bake.cpp +++ b/intern/cycles/render/bake.cpp @@ -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; diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp index 2342dd52d86..01f853dda71 100644 --- a/intern/cycles/render/buffers.cpp +++ b/intern/cycles/render/buffers.cpp @@ -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(); -} - -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(); - } + rgba_byte.free(); + rgba_half.free(); } -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 diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h index 2780fc8a68d..8563d6674ec 100644 --- a/intern/cycles/render/buffers.h +++ b/intern/cycles/render/buffers.h @@ -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 diff --git a/intern/cycles/render/image.cpp b/intern/cycles/render/image.cpp index e7f5ff002b7..625901ff258 100644 --- a/intern/cycles/render/image.cpp +++ b/intern/cycles/render/image.cpp @@ -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; } diff --git a/intern/cycles/render/integrator.cpp b/intern/cycles/render/integrator.cpp index b128f18db08..33c3dac9e81 100644 --- a/intern/cycles/render/integrator.cpp +++ b/intern/cycles/render/integrator.cpp @@ -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) diff --git a/intern/cycles/render/light.cpp b/intern/cycles/render/light.cpp index 9664e1310d5..b3804f34963 100644 --- a/intern/cycles/render/light.cpp +++ b/intern/cycles/render/light.cpp @@ -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*/) diff --git a/intern/cycles/render/mesh.cpp b/intern/cycles/render/mesh.cpp index 685272b80c1..75bdf71616f 100644 --- a/intern/cycles/render/mesh.cpp +++ b/intern/cycles/render/mesh.cpp @@ -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(); diff --git a/intern/cycles/render/mesh_displace.cpp b/intern/cycles/render/mesh_displace.cpp index c06cf86ea9c..ab3ae40d931 100644 --- a/intern/cycles/render/mesh_displace.cpp +++ b/intern/cycles/render/mesh_displace.cpp @@ -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 diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp index daa872239ce..57e44861e40 100644 --- a/intern/cycles/render/object.cpp +++ b/intern/cycles/render/object.cpp @@ -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) diff --git a/intern/cycles/render/particles.cpp b/intern/cycles/render/particles.cpp index a84ca51f274..06ff45b09bd 100644 --- a/intern/cycles/render/particles.cpp +++ b/intern/cycles/render/particles.cpp @@ -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*/) diff --git a/intern/cycles/render/scene.cpp b/intern/cycles/render/scene.cpp index e362a35471d..260a325206c 100644 --- a/intern/cycles/render/scene.cpp +++ b/intern/cycles/render/scene.cpp @@ -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)); } diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp index 4642dcfa9a1..74cfd02e1a4 100644 --- a/intern/cycles/render/session.cpp +++ b/intern/cycles/render/session.cpp @@ -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. */ diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp index a77df55e520..70f6d5bab47 100644 --- a/intern/cycles/render/shader.cpp +++ b/intern/cycles/render/shader.cpp @@ -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) diff --git a/intern/cycles/render/svm.cpp b/intern/cycles/render/svm.cpp index cf0dc97ef3f..db53e366d1e 100644 --- a/intern/cycles/render/svm.cpp +++ b/intern/cycles/render/svm.cpp @@ -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 */ diff --git a/intern/cycles/render/tables.cpp b/intern/cycles/render/tables.cpp index 9d04778abc6..5cda977b7f1 100644 --- a/intern/cycles/render/tables.cpp +++ b/intern/cycles/render/tables.cpp @@ -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) diff --git a/intern/cycles/util/util_vector.h b/intern/cycles/util/util_vector.h index 9e74505b14a..ca6b56c9c7e 100644 --- a/intern/cycles/util/util_vector.h +++ b/intern/cycles/util/util_vector.h @@ -177,6 +177,14 @@ public: } } + T *steal_pointer() + { + T *ptr = data_; + data_ = NULL; + clear(); + return ptr; + } + T* resize(size_t newsize) { if(newsize == 0) { |