[Bf-blender-cvs] [57a0cb797d6] master: Code refactor: avoid some unnecessary device memory copying.
Brecht Van Lommel
noreply at git.blender.org
Sat Oct 21 21:13:47 CEST 2017
Commit: 57a0cb797d60024357a3e3a64c1873844b0178bd
Author: Brecht Van Lommel
Date: Fri Oct 20 04:32:29 2017 +0200
Branches: master
https://developer.blender.org/rB57a0cb797d60024357a3e3a64c1873844b0178bd
Code refactor: avoid some unnecessary device memory copying.
===================================================================
M intern/cycles/device/device_cpu.cpp
M intern/cycles/device/device_cuda.cpp
M intern/cycles/device/device_memory.h
M intern/cycles/device/opencl/opencl_base.cpp
M intern/cycles/device/opencl/opencl_split.cpp
M intern/cycles/render/mesh.cpp
M intern/cycles/render/mesh.h
M intern/cycles/render/scene.h
M intern/cycles/render/svm.cpp
M intern/cycles/render/svm.h
M intern/cycles/render/tables.cpp
M intern/cycles/util/util_vector.h
===================================================================
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,
diff --git a/intern/cycles/render/mesh.cpp b/intern/cycles/render/mesh.cpp
index efec1d3e491..69c21fc3cb3 100644
--- a/intern/cycles/render/mesh.cpp
+++ b/intern/cycles/render/mesh.cpp
@@ -1127,14 +1127,12 @@ bool Mesh::is_instanced() const
MeshManager::MeshManager()
{
- bvh = NULL;
need_update = true;
need_flags_update = true;
}
MeshManager::~MeshManager()
{
- delete bvh;
}
void MeshManager::update_osl_attributes(Device *device, Scene *scene, vector<AttributeRequestSet>& mesh_attributes)
@@ -1393,11 +1391,11 @@ static void update_attribute_element_size(Mesh *mesh,
}
static void update_attribute_element_offset(Mesh *mesh,
- vector<float>& attr_float,
+ device_vector<float>& attr_float,
size_t& attr_float_offset,
- vector<float4>& attr_float3,
+ device_vector<float4>& attr_float3,
size_t& attr_float3_offset,
- vector<uchar4>& attr_uchar4,
+ device_vector<uchar4>& attr_uchar4,
size_t& attr_uchar4_offset,
Attribute *mattr,
AttributePrimitive prim,
@@ -1425,7 +1423,7 @@ static void update_attribute_element_offset(Mesh *mesh,
uchar4 *data = mattr->data_uchar4();
offset = attr_uchar4_offset;
- assert(attr_uchar4.capacity() >= offset + size);
+ assert(attr_uchar4.size() >= offset + size);
for(size_t k = 0; k < size; k++) {
attr_uchar4[offset+k] = data[k];
}
@@ -1435,7 +1433,7 @@ static void update_attribute_element_offset(Mesh *mesh,
float *data = mattr->data_float();
offset = attr_float_offset;
- assert(attr_float.capacity() >= offset + size);
+ assert(attr_float.size() >= offset + size);
for(size_t k = 0; k < size; k++) {
attr_float[offset+k] = data[k];
}
@@ -1445,7 +1443,7 @@ static void update_attribute_element_offset(Mesh *mesh,
Transform *tfm = mattr->data_transform();
offset = attr_float3_offset;
- assert(attr_float3.capacity() >= offset + size * 4);
+ assert(attr_float3.size() >= offset + size * 4);
for(size_t k = 0; k < size*4; k++) {
attr_float3[offset+k] = (&tfm->x)[k];
}
@@ -1455,7 +1453,7 @@ static void update_attribute_element_offset(Mesh *mesh,
float4 *data = mattr->data_float4();
offset = attr_float3_offset;
- assert(attr_float3.capacity() >= offset + size);
+ assert(attr_float3.size() >= offset + size);
for(size_t k = 0; k < size; k++) {
attr_float3[offset+k] = data[k];
}
@@ -1556,9 +1554,9 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
}
}
- vector<float> attr_float(attr_float_size);
- vector<float4> attr_float3(attr_float3_size);
- vector<uchar4> attr_uchar4(attr_uchar4_size);
+ dscene->attributes_float.resize(attr_float_size);
+ dscene->attributes_float3.resize(attr_float3_size);
+ dscene->attributes_uchar4.resize(attr_uchar4_size);
size_t attr_float_offset = 0;
size_t attr_float3_offset = 0;
@@ -1577,27 +1575,27 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
Attribute *subd_mattr = mesh->subd_attributes.find(req);
update_attribute_element_offset(mesh,
- attr_float, attr_float_offset,
- attr_float3, attr_float3_offset,
- attr_uchar4, attr_uchar4_offset,
+ dscene->attributes_float, attr_float_offset,
+ dscene->attributes_float3, attr_float3_offset,
+ dscene->attributes_uchar4, attr_uchar4_offset,
triangle_mattr,
ATTR_PRIM_TRIANGLE,
req.triangle_type,
req.triangle_desc);
update_attribute_element_offset(mesh,
- attr_float, attr_float_offset,
- attr_float3, attr_float3_offset,
- attr_uchar4, attr_uchar4_offset,
+ dscene->attributes_float, attr_float_offset,
+ dscene->attributes_float3, attr_float3_offset,
+ dscene->attributes_uchar4, attr_uchar4_offset,
curve_mattr,
ATTR_PRIM_CURVE,
req.curve_type,
req.curve_desc);
update_attribute_element_offset(mesh,
- attr_float, attr_float_offset,
- attr_float3, attr_float3_offset,
- attr_uchar4, attr_uchar4_offset,
+ dscene->attributes_float, attr_float_offset,
+ dscene->attributes_float3, attr_float3_offset,
+ dscene->attributes_uchar4, attr_uchar4_offset,
subd_mattr,
ATTR_PRIM_SUBD,
req.subd_type,
@@ -1618,16 +1616,13 @@ void MeshManager::device_update_attributes(Device *device, De
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list