Code refactor: move more memory allocation logic into device API.
authorBrecht Van Lommel <brechtvanlommel@gmail.com>
Fri, 20 Oct 2017 23:09:59 +0000 (01:09 +0200)
committerBrecht Van Lommel <brechtvanlommel@gmail.com>
Mon, 23 Oct 2017 23:25:19 +0000 (01:25 +0200)
* Remove tex_* and pixels_* functions, replace by mem_*.
* Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices.
* No longer create device_memory and call mem_* directly, always go
  through device_only_memory, device_vector and device_pixels.

30 files changed:
intern/cycles/device/device.cpp
intern/cycles/device/device.h
intern/cycles/device/device_cpu.cpp
intern/cycles/device/device_cuda.cpp
intern/cycles/device/device_denoising.cpp
intern/cycles/device/device_memory.cpp
intern/cycles/device/device_memory.h
intern/cycles/device/device_multi.cpp
intern/cycles/device/device_network.cpp
intern/cycles/device/device_network.h
intern/cycles/device/device_split_kernel.cpp
intern/cycles/device/opencl/memory_manager.cpp
intern/cycles/device/opencl/opencl_base.cpp
intern/cycles/device/opencl/opencl_split.cpp
intern/cycles/render/bake.cpp
intern/cycles/render/buffers.cpp
intern/cycles/render/buffers.h
intern/cycles/render/image.cpp
intern/cycles/render/integrator.cpp
intern/cycles/render/light.cpp
intern/cycles/render/mesh.cpp
intern/cycles/render/mesh_displace.cpp
intern/cycles/render/object.cpp
intern/cycles/render/particles.cpp
intern/cycles/render/scene.cpp
intern/cycles/render/session.cpp
intern/cycles/render/shader.cpp
intern/cycles/render/svm.cpp
intern/cycles/render/tables.cpp
intern/cycles/util/util_vector.h

index 9de10c184fb60409cb4778e8220b6626d03966a4..41fbe7ce81bc92ce711fdb0874db2bdf092069cb 100644 (file)
@@ -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);
index 6bb65cde2a38d9f26ad7b2534046ae6717a57429..316bf70a5c3cb782d84069be2fb1dbf3516f5656 100644 (file)
@@ -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;
index b4398f210144ed533eb5c27ace72950a63421dbe..32ab18fe1646e32b5ba67bbfef2e270fce49ac3d 100644 (file)
@@ -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;
        }
 
index be606a92434b3ad9f8491d92de2b66d980ff66ca..aa6386e455b176dcadad4b60ac6c7c22147ad05a 100644 (file)
@@ -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;
 }
index 2c3bfefd8b08604535c6d0b0c93ec5c97a186186..2d39721e3d338b74877ec5d9c34befb53d1a316f 100644 (file)
@@ -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;
 }
 
index 98fa638ef8e8f713bb4208fb958d2dcd9f99c700..9f4f60e7531bf4a758ef165d70d73f1af4b28add 100644 (file)
@@ -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)
index 3dfecde59d8ecd065256749964499180391a3b90..7bf8bdc1ceac5eaa91850f5d8ad68c9bcd863641 100644 (file)
 
 /* 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;
 };
index 7f7fbc0d1d32f9060cbd632be3e2708dd6137224..0a6dd90c86d1d0cc055bcc4aee6e987425dc46f2 100644 (file)
@@ -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();
                                }
                        }
                }
index bdc88b6acae280bc89507f98208b46709e1894dd..fa231c817e604cb6cc67346309a7bb0127eea204 100644 (file)
@@ -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);
 
index 8a53290f421e2f9f6c762bf79ea9556503bccba5..a38d962c0af9ab83d0740088a32c89ffe038fb3c 100644 (file)
@@ -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)
index 6c8befa89bead042d6ccc380e46f12831729446a..f2839a8b1b9ecc6d793ed1b11b39424b3216b3fc 100644 (file)
@@ -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;
 
index e48367b8987d1189a20d1fdd87737715f19c1f09..a791b374774e7962e8de5a342bdd5c5415c10006 100644 (file)
@@ -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()
index 90f461b4c988269e26b14fe9ec41bc89a94c296e..5e9debc3b17bd1af4804f15297571c73974caddb 100644 (file)
@@ -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);
 
index c966ebe0c5ea403a21bd1b22ba15bd49abaa8d26..96139afa45022fc357a1c37534734f649753d23d 100644 (file)
@@ -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()",
index 66615bf336c2019639cfe523b8e9b7a078ec8755..99f68b6aa007aa818df75fda86f35c25d37db4a0 100644 (file)
@@ -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;
index 2342dd52d866099efb43bdb9998f8fae20440610..01f853dda718621b841179b7e81a6bc6e9cd2a9c 100644 (file)
@@ -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
 
index 2780fc8a68df4fe1e69dcdd4cc9f84d850bd6856..8563d6674ec398d6cfabd697be11d1d48535e30f 100644 (file)
@@ -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
index e7f5ff002b7bb80df08457d6ca414b84fe9351f3..625901ff258ea33acb72f0e01ff50a9f8fc2d8a4 100644 (file)
@@ -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;
                }
 
index b128f18db081605a1b1e9da9bfb241638e7e721f..33c3dac9e81560568a1d1965cd992b90edee502c 100644 (file)
@@ -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)
index 9664e1310d5e7d1f57b08bd13a14a6a7729e28ac..b3804f34963ff284085ca21b29e0af57a42545f3 100644 (file)
@@ -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*/)
index 685272b80c12dc6da8548a04174a5874b02ae0e0..75bdf71616f1352cbec8fe250d294971f60af3a6 100644 (file)
@@ -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();
index c06cf86ea9cca709b110a179ac656e0f992a23de..ab3ae40d931175bc4b07007f1d95ad6401876c0f 100644 (file)
@@ -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
index daa872239ce69f63174cb5043edb351840d94199..57e44861e4054d539ad1bd8e8ca3899bdea30b70 100644 (file)
@@ -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)
index a84ca51f2749cfd4a84983f09ca9c93999b21ace..06ff45b09bde51c59820c50b7cca7d54552e4ab4 100644 (file)
@@ -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*/)
index e362a35471dd06100858031b17440ea7e811b7e3..260a325206c847dc2fd72cc0bc738a87a41a546b 100644 (file)
 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));
 }
index 4642dcfa9a1297a80d03df2e62cddaf5c0bb7b70..74cfd02e1a46359000884554e188911f2eb7b2e1 100644 (file)
@@ -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. */
index a77df55e52022e4a648d379f66f0748f1eeae848..70f6d5bab474ab4c5a3f55d0dd0692422c8c5bb2 100644 (file)
@@ -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)
index cf0dc97ef3fb90279da4009638196aecad7fc2bb..db53e366d1ecbe2c986b10ef95eb47bd2f629631 100644 (file)
@@ -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 */
index 9d04778abc68e2eb68636663ea37070f7caccdfa..5cda977b7f1b53e0f537428057016bbd38bcc992 100644 (file)
@@ -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)
index 9e74505b14abab44a328d3263952503fc5b88b26..ca6b56c9c7e3ba9c9b748d1464fc6c9adb7e05b3 100644 (file)
@@ -177,6 +177,14 @@ public:
                }
        }
 
+       T *steal_pointer()
+       {
+               T *ptr = data_;
+               data_ = NULL;
+               clear();
+               return ptr;
+       }
+
        T* resize(size_t newsize)
        {
                if(newsize == 0) {