Skip to content

Commit

Permalink
Bugfix agent death (#127)
Browse files Browse the repository at this point in the history 
* BugFix: Render the number of elements in the buffer

Not, the number eleements at next update if less.

* Triple buffered CUDATextureBuffer

This removes undefined behaviour.

Closes #126
  • Loading branch information
@Robadob
Robadob authored Mar 31, 2023
1 parent 02097b5 commit dea46cd
Show file tree
Hide file tree
Showing 3 changed files with 61 additions and 15 deletions.
38 changes: 28 additions & 10 deletions src/flamegpu/visualiser/Visualiser.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -528,15 +528,17 @@ void Visualiser::renderAgentStates() {
// Alloc new buffs (this needs to occur in render thread!)
tb = mallocGLInteropTextureBuffer<float>(newSize * TexBufferConfig::SamplerElements(_tb.first), 1);
// Copy any old data to the buffer
if (old_tb && tb && old_tb->d_mappedPointer && tb->d_mappedPointer && as.dataSize)
_cudaMemcpyDeviceToDevice(tb->d_mappedPointer, old_tb->d_mappedPointer, as.dataSize * sizeof(float) * TexBufferConfig::SamplerElements(_tb.first));
if (old_tb && tb && old_tb->d_pointer && tb->d_pointer && as.dataSize) {
_cudaMemcpyDeviceToDevice(tb->d_pointer, old_tb->d_pointer, as.dataSize * sizeof(float) * TexBufferConfig::SamplerElements(_tb.first));
tb->d_pointer_outofdate = true;
}
// Bind texture name to texture unit
GL_CALL(glActiveTexture(GL_TEXTURE0 + as.tex_unit_offset + tui));
GL_CALL(glBindTexture(GL_TEXTURE_BUFFER, tb->glTexName));
GL_CALL(glActiveTexture(GL_TEXTURE0));
shader_vec->addTexture(samplerName.c_str(), GL_TEXTURE_BUFFER, tb->glTexName, as.tex_unit_offset + tui);
// Free old buff
if (old_tb && old_tb->d_mappedPointer)
if (old_tb && old_tb->d_pointer)
freeGLInteropTextureBuffer(old_tb);
++tui;
GL_CHECK();
Expand All @@ -549,15 +551,17 @@ void Visualiser::renderAgentStates() {
// Alloc new buffs (this needs to occur in other thread!!!)
tb = mallocGLInteropTextureBuffer<float>(newSize * _tb.second.first.array_length * TexBufferConfig::SamplerElements(_tb.first), 1);
// Copy any old data to the buffer
if (old_tb && tb && old_tb->d_mappedPointer && tb->d_mappedPointer && as.dataSize)
_cudaMemcpyDeviceToDevice(tb->d_mappedPointer, old_tb->d_mappedPointer, as.dataSize * _tb.second.first.array_length * sizeof(float) * TexBufferConfig::SamplerElements(_tb.first));
if (old_tb && tb && old_tb->d_pointer && tb->d_pointer && as.dataSize) {
_cudaMemcpyDeviceToDevice(tb->d_pointer, old_tb->d_pointer, as.dataSize * _tb.second.first.array_length * sizeof(float) * TexBufferConfig::SamplerElements(_tb.first));
tb->d_pointer_outofdate = true;
}
// Bind texture name to texture unit
GL_CALL(glActiveTexture(GL_TEXTURE0 + as.tex_unit_offset + tui));
GL_CALL(glBindTexture(GL_TEXTURE_BUFFER, tb->glTexName));
GL_CALL(glActiveTexture(GL_TEXTURE0));
shader_vec->addTexture(_tb.second.first.nameInShader.c_str(), GL_TEXTURE_BUFFER, tb->glTexName, as.tex_unit_offset + tui);
// Free old buff
if (old_tb && old_tb->d_mappedPointer)
if (old_tb && old_tb->d_pointer)
freeGLInteropTextureBuffer(old_tb);
++tui;
GL_CHECK();
Expand Down Expand Up @@ -591,10 +595,22 @@ void Visualiser::renderAgentStates() {

// Render agents
if (closeSplashScreen) {
// Update data in mapped buffers
for (auto& _as : agentStates) {
auto& as = _as.second;
if (as.core_texture_buffers.begin()->second) {
for (auto& _tb : as.core_texture_buffers) {
_tb.second->updateMapped();
}
for (auto& _tb : as.custom_texture_buffers) {
_tb.second.second->updateMapped();
}
}
}
GL_CHECK();
for (auto &as : agentStates) {
if (!as.second.core_texture_buffers.empty() && as.second.dataSize) // Check to make sure buffer has been allocated successfully
if (as.second.requiredSize) // Also check we actually have agents (buffer might be bigger than the agents)
as.second.entity->renderInstances(static_cast<int>(std::min(as.second.dataSize, as.second.requiredSize)));
as.second.entity->renderInstances(static_cast<int>(as.second.dataSize));
}
}
if (guard)
Expand Down Expand Up @@ -625,13 +641,15 @@ void Visualiser::updateAgentStateBuffer(const std::string &agent_name, const std
for (const auto &_ext_tb : ext_core_tex_buffers) {
auto &ext_tb = _ext_tb.second;
auto &int_tb = as.core_texture_buffers.at(_ext_tb.first);
visassert(_cudaMemcpyDeviceToDevice(int_tb->d_mappedPointer, ext_tb.t_d_ptr, as.dataSize * sizeof(float) * TexBufferConfig::SamplerElements(_ext_tb.first)));
visassert(_cudaMemcpyDeviceToDevice(int_tb->d_pointer, ext_tb.t_d_ptr, as.dataSize * sizeof(float) * TexBufferConfig::SamplerElements(_ext_tb.first)));
int_tb->d_pointer_outofdate = true;
}
for (const auto& _ext_tb : ext_tex_buffers) {
auto& ext_tb = _ext_tb.second;
for (auto int_tb = as.custom_texture_buffers.find(_ext_tb.first); int_tb != as.custom_texture_buffers.end(); ++int_tb) {
if (ext_tb.nameInShader == int_tb->second.first.nameInShader) {
visassert(_cudaMemcpyDeviceToDevice(int_tb->second.second->d_mappedPointer, ext_tb.t_d_ptr, as.dataSize * ext_tb.array_length * sizeof(float) * TexBufferConfig::SamplerElements(_ext_tb.first)));
visassert(_cudaMemcpyDeviceToDevice(int_tb->second.second->d_pointer, ext_tb.t_d_ptr, as.dataSize * ext_tb.array_length * sizeof(float) * TexBufferConfig::SamplerElements(_ext_tb.first)));
int_tb->second.second->d_pointer_outofdate = true;
}
}
}
Expand Down
16 changes: 16 additions & 0 deletions src/flamegpu/visualiser/util/cuda.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -156,6 +156,19 @@ void freeGLInteropTextureBuffer(CUDATextureBuffer<T> *texBuf) {
GL_CALL(glDeleteTextures(1, &texBuf->glTexName));
delete texBuf;
}

template<class T>
bool CUDATextureBuffer<T>::updateMapped() {
if (!this->d_pointer_outofdate) return true;
CUDA_CALL(cudaGraphicsMapResources(1, const_cast<cudaGraphicsResource_t *>(&this->cuGraphicsRes)));
void* d_mappedPointer = nullptr;
CUDA_CALL(cudaGraphicsResourceGetMappedPointer(&d_mappedPointer, 0, this->cuGraphicsRes));
auto t = cudaMemcpy(d_mappedPointer, this->d_pointer, this->elementCount * this->componentCount * sizeof(T), cudaMemcpyDeviceToDevice);
CUDA_CALL(cudaGraphicsUnmapResources(1, const_cast<cudaGraphicsResource_t*>(&this->cuGraphicsRes), 0));
CUDA_CALL(t);
this->d_pointer_outofdate = false;
return t == cudaSuccess;
}
/**
* Returns true if cudaMemcpy returns cudaSuccess
*/
Expand All @@ -171,6 +184,9 @@ template CUDATextureBuffer<unsigned int> *mallocGLInteropTextureBuffer(const uns
template void freeGLInteropTextureBuffer(CUDATextureBuffer<float>*);
template void freeGLInteropTextureBuffer(CUDATextureBuffer<int>*);
template void freeGLInteropTextureBuffer(CUDATextureBuffer<unsigned int>*);
template bool CUDATextureBuffer<float>::updateMapped();
template bool CUDATextureBuffer<int>::updateMapped();
template bool CUDATextureBuffer<unsigned int>::updateMapped();

} // namespace visualiser
} // namespace flamegpu
22 changes: 17 additions & 5 deletions src/flamegpu/visualiser/util/cuda.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,31 +17,43 @@ typedef uint64_t cudaTextureObject_t; // actual header says unsigned long long,
#include <cuda_runtime.h>
#endif

/**
* Represents a double buffered CUDATextureBuffer
* Copy data into d_pointer
* Set d_pointer_outofdate = true
* Then call updateMapped() before using glTexName/glTBO
*/
template<class T>
struct CUDATextureBuffer {
CUDATextureBuffer(
const GLuint glTexName,
const GLuint glTBO,
T *d_mappedPointer,
const cudaGraphicsResource_t cuGraphicsRes,
cudaGraphicsResource_t cuGraphicsRes,
cudaTextureObject_t cuTextureObj,
const unsigned int elementCount,
const unsigned int componentCount
)
: glTexName(glTexName)
, glTBO(glTBO)
, d_mappedPointer(d_mappedPointer)
, d_pointer(d_mappedPointer)
, cuGraphicsRes(cuGraphicsRes)
, cuTextureObj(cuTextureObj)
, elementCount(elementCount)
, componentCount(componentCount) { }
const GLuint glTexName;
const GLuint glTBO;
T *d_mappedPointer;
const cudaGraphicsResource_t cuGraphicsRes; // These are typedefs over pointers, need to store the actual struct?
const cudaTextureObject_t cuTextureObj;
T *d_pointer;
bool d_pointer_outofdate = false;
const cudaGraphicsResource_t cuGraphicsRes;
const cudaTextureObject_t cuTextureObj; // Using this is currently unsafe, the cuGraphicRes should technically be mapped/unmapped around use
const unsigned int elementCount;
const unsigned int componentCount;
/**
* Copy data from d_pointer to glTBO
* @return true on cudaSuccess
*/
bool updateMapped();
};
/**
* Allocates a GL_TEXTURE_BUFFER of the desired size and binds it for use with CUDA-GL interop
Expand Down

0 comments on commit dea46cd

Please sign in to comment.