diff options
author | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2017-10-20 05:32:29 +0300 |
---|---|---|
committer | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2017-10-21 21:58:28 +0300 |
commit | 57a0cb797d60024357a3e3a64c1873844b0178bd (patch) | |
tree | 1f8ade576fbbc6cbbf7a41c51304ee8ae3fe95b6 /intern/cycles/device | |
parent | 92ec4863c22f249a21a5b5224d91fcab5c602100 (diff) |
Code refactor: avoid some unnecessary device memory copying.
Diffstat (limited to 'intern/cycles/device')
-rw-r--r-- | intern/cycles/device/device_cpu.cpp | 2 | ||||
-rw-r--r-- | intern/cycles/device/device_cuda.cpp | 5 | ||||
-rw-r--r-- | intern/cycles/device/device_memory.h | 36 | ||||
-rw-r--r-- | intern/cycles/device/opencl/opencl_base.cpp | 13 | ||||
-rw-r--r-- | intern/cycles/device/opencl/opencl_split.cpp | 3 |
5 files changed, 25 insertions, 34 deletions
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp index af1bbc0db18..0ba00da16a6 100644 --- a/intern/cycles/device/device_cpu.cpp +++ b/intern/cycles/device/device_cpu.cpp @@ -379,7 +379,7 @@ public: texture_info.resize(flat_slot + 128); } - TextureInfo& info = texture_info.get_data()[flat_slot]; + TextureInfo& info = texture_info[flat_slot]; info.data = (uint64_t)mem.data_pointer; info.cl_buffer = 0; info.interpolation = interpolation; diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp index c245d7d8408..0f17b67c8c6 100644 --- a/intern/cycles/device/device_cuda.cpp +++ b/intern/cycles/device/device_cuda.cpp @@ -840,7 +840,7 @@ public: } /* Set Mapping and tag that we need to (re-)upload to device */ - TextureInfo& info = texture_info.get_data()[flat_slot]; + TextureInfo& info = texture_info[flat_slot]; info.data = (uint64_t)tex; info.cl_buffer = 0; info.interpolation = interpolation; @@ -1911,9 +1911,10 @@ uint64_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory 0, 0, (void**)&args, 0)); device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t)); + size_t size = size_buffer[0]; device->mem_free(size_buffer); - return *size_buffer.get_data(); + return size; } bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim, diff --git a/intern/cycles/device/device_memory.h b/intern/cycles/device/device_memory.h index 20707ad04c9..eeeca61496e 100644 --- a/intern/cycles/device/device_memory.h +++ b/intern/cycles/device/device_memory.h @@ -270,31 +270,14 @@ public: return &data[0]; } - T *copy(T *ptr, size_t width, size_t height = 0, size_t depth = 0) + void steal_data(array<T>& from) { - T *mem = resize(width, height, depth); - if(mem != NULL) { - memcpy(mem, ptr, memory_size()); - } - return mem; - } - - void copy_at(T *ptr, size_t offset, size_t size) - { - if(size > 0) { - size_t mem_size = size*data_elements*datatype_size(data_type); - memcpy(&data[0] + offset, ptr, mem_size); - } - } - - void reference(T *ptr, size_t width, size_t height = 0, size_t depth = 0) - { - data.clear(); - data_size = width * ((height == 0)? 1: height) * ((depth == 0)? 1: depth); - data_pointer = (device_ptr)ptr; - data_width = width; - data_height = height; - data_depth = depth; + data.steal_data(from); + data_size = data.size(); + data_pointer = (data_size)? (device_ptr)&data[0]: 0; + data_width = data_size; + data_height = 0; + data_depth = 0; } void clear() @@ -318,6 +301,11 @@ public: return &data[0]; } + T& operator[](size_t i) + { + return data[i]; + } + private: array<T> data; }; diff --git a/intern/cycles/device/opencl/opencl_base.cpp b/intern/cycles/device/opencl/opencl_base.cpp index 6747a8a83ac..48c32a9dc5c 100644 --- a/intern/cycles/device/opencl/opencl_base.cpp +++ b/intern/cycles/device/opencl/opencl_base.cpp @@ -494,20 +494,21 @@ void OpenCLDeviceBase::mem_free_sub_ptr(device_ptr device_pointer) void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size) { ConstMemMap::iterator i = const_mem_map.find(name); + device_vector<uchar> *data; if(i == const_mem_map.end()) { - device_vector<uchar> *data = new device_vector<uchar>(); - data->copy((uchar*)host, size); + data = new device_vector<uchar>(); + data->resize(size); mem_alloc(name, *data, MEM_READ_ONLY); - i = const_mem_map.insert(ConstMemMap::value_type(name, data)).first; + const_mem_map.insert(ConstMemMap::value_type(name, data)); } else { - device_vector<uchar> *data = i->second; - data->copy((uchar*)host, size); + data = i->second; } - mem_copy_to(*i->second); + memcpy(data->get_data(), host, size); + mem_copy_to(*data); } void OpenCLDeviceBase::tex_alloc(const char *name, diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp index b4e9419ebbd..920106f92d4 100644 --- a/intern/cycles/device/opencl/opencl_split.cpp +++ b/intern/cycles/device/opencl/opencl_split.cpp @@ -309,6 +309,7 @@ public: device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel"); device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint64_t)); + size_t size = size_buffer[0]; device->mem_free(size_buffer); if(device->ciErr != CL_SUCCESS) { @@ -318,7 +319,7 @@ public: return 0; } - return *size_buffer.get_data(); + return size; } virtual bool enqueue_split_kernel_data_init(const KernelDimensions& dim, |