From 3b4f85f1017a09cdb5cdc2c656ec140f595d38a8 Mon Sep 17 00:00:00 2001 From: Sergen Eren Date: Fri, 27 Dec 2019 16:46:19 +0300 Subject: [PATCH] will allocate and fill procedural volume data in texture_kernels.cu --- source/gpu_vdb/gpu_vdb.cpp | 82 ++++++++++++++++++++++++++++++++++++++ source/gpu_vdb/gpu_vdb.h | 24 +++++++++++ source/main.cpp | 2 - source/texture_kernels.cu | 15 ++++++- 4 files changed, 120 insertions(+), 3 deletions(-) diff --git a/source/gpu_vdb/gpu_vdb.cpp b/source/gpu_vdb/gpu_vdb.cpp index 5ed8411..b63f6f3 100644 --- a/source/gpu_vdb/gpu_vdb.cpp +++ b/source/gpu_vdb/gpu_vdb.cpp @@ -470,3 +470,85 @@ bool GPU_VDB::loadVDB(std::string filename, std::string density_channel, std::st set_xform(xform_temp); return true; } + + +// Class implementations for procedural volume + + +GPU_PROC_VOL::~GPU_PROC_VOL() { + + + if (device_density_buffer) { + cudaFree(device_density_buffer); + } + +} + + +GPU_PROC_VOL::GPU_PROC_VOL(const GPU_PROC_VOL& copy){ + + set_xform(copy.get_xform()); + this->vdb_info = copy.vdb_info; + +} + +GPU_PROC_VOL::GPU_PROC_VOL() { + + CUresult error = cuModuleLoad(&texture_module, "texture_kernels.ptx"); + if (error != CUDA_SUCCESS) log("cuModuleLoad" + std::to_string(error), ERROR); + + error = cuModuleGetFunction(&fill_buffer_function, texture_module, "fill_volume_buffer"); + if (error != CUDA_SUCCESS) { + log("Unable to bind buffer fill function!", ERROR); + } + +} + +// fill vdb_info density texture with procedural noise texture +bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) { + + if (min.x > max.x&& min.y > max.y&& min.z > max.z) { + log("max < min", ERROR); + return false; + } + + mat4 xform; + xform.scale(make_float3(res)); + set_xform(xform); + + int dim_x = floorf((max.x - min.x) / res); + int dim_y = floorf((max.y - min.y) / res); + int dim_z = floorf((max.z - min.z) / res); + + dimensions = make_int3(dim_x, dim_y, dim_z); + + // Fill vdb info parameters that would normally come from a vdb file + vdb_info.dim = dimensions; + vdb_info.bmin = make_float3(.0f); + vdb_info.bmax = make_float3(dimensions); + vdb_info.voxelsize = res; + vdb_info.min_density = .0f; + vdb_info.max_density = 1.0f; + vdb_info.has_emission = false; + vdb_info.has_color = false; + + // set noise type , see texture_kernels.cu for noise types + int noise_type = 0; + + // Allocate device memory for volume buffer + checkCudaErrors(cudaMalloc(&device_density_buffer, dimensions.x * dimensions.y * dimensions.z * sizeof(float))); + + dim3 block(16, 16, 16); + dim3 grid(int(dimensions.x / block.x) + 1, int(dimensions.y / block.y) + 1, int(dimensions.z / block.z) + 1); + + 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 + + + + + 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 33384b7..d65545f 100644 --- a/source/gpu_vdb/gpu_vdb.h +++ b/source/gpu_vdb/gpu_vdb.h @@ -39,6 +39,7 @@ #ifndef _GPU_VDB_H_ #define _GPU_VDB_H_ +#include "cuda.h" #include "cuda_runtime_api.h" #include "texture_types.h" @@ -145,10 +146,33 @@ class GPU_VDB { } VDB_INFO vdb_info; + private: mat4 xform; }; + +class GPU_PROC_VOL : virtual public GPU_VDB { + + +public: + + __host__ GPU_PROC_VOL(); + __host__ GPU_PROC_VOL(const GPU_PROC_VOL& copy); + __host__ ~GPU_PROC_VOL(); + + __host__ bool create_volume(float3 min, float3 max, float res); + +private: + + CUmodule texture_module; + CUfunction fill_buffer_function; + float *device_density_buffer; + float resolution; + int3 dimensions; +}; + + #endif //endif _GPU_VDB_H_ \ No newline at end of file diff --git a/source/main.cpp b/source/main.cpp index 26cc360..084df52 100644 --- a/source/main.cpp +++ b/source/main.cpp @@ -1846,8 +1846,6 @@ int main(const int argc, const char* argv[]) // Launch volume rendering kernel. dim3 block(16, 16, 1); dim3 grid(int(width / block.x) + 1, int(height / block.y) + 1, 1); - dim3 threads_per_block(16, 16); - dim3 num_blocks((width + 15) / 16, (height + 15) / 16); void *params[] = { &cam, (void *)&l_list , (void *)&d_volume_ptr, (void *)&d_geo_ptr, (void*)&d_geo_list_ptr ,&bvh_builder.bvh.BVHNodes, &bvh_builder.root ,(void *)atmos_params, &kernel_params}; cuLaunchKernel(cuRaycastKernel, grid.x, grid.y, 1, block.x, block.y, 1, 0, NULL, params, NULL); diff --git a/source/texture_kernels.cu b/source/texture_kernels.cu index 6cc0b84..74e2c99 100644 --- a/source/texture_kernels.cu +++ b/source/texture_kernels.cu @@ -58,7 +58,6 @@ extern "C" __global__ void glow(const Kernel_params kernel_params, float treshold , const int width, const int height) { - int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x >= width || y >= height) return; @@ -67,4 +66,18 @@ extern "C" __global__ void glow(const Kernel_params kernel_params, float treshol // TODO gaussian blur and add glow effect to display buffer +} + +extern "C" __global__ void fill_volume_buffer(float *buffer, int3 dims, int noise_type) { + + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + int z = blockIdx.z * blockDim.z + threadIdx.z; + + if (x >= dims.x || y >= dims.y || z >= dims.z) return; + + const unsigned int idx = x + dims.x * (y + dims.y * z); + + buffer[idx] = .0f; + } \ No newline at end of file