From 51972d84146daa026d85e610a39156f6eefb127f Mon Sep 17 00:00:00 2001 From: Sergen Eren Date: Fri, 27 Dec 2019 17:06:10 +0300 Subject: [PATCH] fill volume buffer in gpu and send it to vdb_info density texture --- source/gpu_vdb/gpu_vdb.cpp | 44 +++++++++++++++++++++++++++++++++++++- source/gpu_vdb/gpu_vdb.h | 2 +- 2 files changed, 44 insertions(+), 2 deletions(-) diff --git a/source/gpu_vdb/gpu_vdb.cpp b/source/gpu_vdb/gpu_vdb.cpp index b63f6f3..5cb72cc 100644 --- a/source/gpu_vdb/gpu_vdb.cpp +++ b/source/gpu_vdb/gpu_vdb.cpp @@ -544,11 +544,53 @@ bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) { void* params[] = {&device_density_buffer, (void *)&dimensions, &noise_type}; cuLaunchKernel(fill_buffer_function, grid.x, grid.y, grid.z, block.x, block.y, block.z, 0, NULL, params, NULL); - // TODO send buffer to texture + // send buffer to texture + cudaExtent vol_size; + vol_size.width = dim_x; + vol_size.height = dim_y; + vol_size.depth = dim_z; + float* volume_data_host = (float*)malloc(dim_x * dim_y * dim_z * sizeof(float)); + checkCudaErrors(cudaMemcpy(volume_data_host, device_density_buffer, dim_x * dim_y * dim_z * sizeof(float), cudaMemcpyDeviceToHost)); + // create 3D array + cudaArray* d_volumeArray = 0; + + cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); + checkCudaErrors(cudaMalloc3DArray(&d_volumeArray, &channelDesc, vol_size)); + + // copy data to 3D array + cudaMemcpy3DParms copyParams = { 0 }; + copyParams.srcPtr = make_cudaPitchedPtr(volume_data_host, vol_size.width * sizeof(float), vol_size.width, vol_size.height); + copyParams.dstArray = d_volumeArray; + copyParams.extent = vol_size; + copyParams.kind = cudaMemcpyHostToDevice; + checkCudaErrors(cudaMemcpy3D(©Params)); + + + cudaResourceDesc texRes; + memset(&texRes, 0, sizeof(cudaResourceDesc)); + + texRes.resType = cudaResourceTypeArray; + texRes.res.array.array = d_volumeArray; + + cudaTextureDesc texDescr; + memset(&texDescr, 0, sizeof(cudaTextureDesc)); + + texDescr.normalizedCoords = true; // access with normalized texture coordinates + texDescr.filterMode = cudaFilterModeLinear; // linear interpolation + + texDescr.addressMode[0] = cudaAddressModeClamp; // clamp texture coordinates + texDescr.addressMode[1] = cudaAddressModeClamp; + texDescr.addressMode[2] = cudaAddressModeClamp; + + texDescr.readMode = cudaReadModeElementType; + //texDescr.readMode = cudaReadModeNormalizedFloat; + + checkCudaErrors(cudaCreateTextureObject(&vdb_info.density_texture, &texRes, &texDescr, NULL)); cudaFree(device_density_buffer); + return true; } \ No newline at end of file diff --git a/source/gpu_vdb/gpu_vdb.h b/source/gpu_vdb/gpu_vdb.h index d65545f..d9ce45c 100644 --- a/source/gpu_vdb/gpu_vdb.h +++ b/source/gpu_vdb/gpu_vdb.h @@ -154,7 +154,7 @@ class GPU_VDB { }; -class GPU_PROC_VOL : virtual public GPU_VDB { +class GPU_PROC_VOL : public GPU_VDB { public: