From c3f853d5e6a8886c4108687a0504b6cf5149d228 Mon Sep 17 00:00:00 2001 From: Sergen Eren Date: Fri, 27 Dec 2019 19:18:52 +0300 Subject: [PATCH] pass the kernel functions in class function. Test volume with a new gui --- source/gpu_vdb/gpu_vdb.cpp | 27 ++------------------ source/gpu_vdb/gpu_vdb.h | 2 +- source/main.cpp | 39 +++++++++++++++++++++++++++-- source/texture_kernels.cu | 50 ++++++++++++++++++++++++++++++++++++-- 4 files changed, 88 insertions(+), 30 deletions(-) diff --git a/source/gpu_vdb/gpu_vdb.cpp b/source/gpu_vdb/gpu_vdb.cpp index 90659dd..5b6d6ee 100644 --- a/source/gpu_vdb/gpu_vdb.cpp +++ b/source/gpu_vdb/gpu_vdb.cpp @@ -505,7 +505,7 @@ GPU_PROC_VOL::GPU_PROC_VOL() { } // fill vdb_info density texture with procedural noise texture -bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) { +bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res, int noise_type, float scale) { log("Creating procedural volume...", LOG); @@ -539,10 +539,6 @@ bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) { 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 log("Allocating device memory for volume buffer...", LOG); @@ -552,7 +548,7 @@ bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) { dim3 grid(int(dimensions.x / block.x) + 1, int(dimensions.y / block.y) + 1, int(dimensions.z / block.z) + 1); log("filling volume buffer in device...", LOG); - void* params[] = { &device_density_buffer , &dimensions, &noise_type}; + void* params[] = { &device_density_buffer , &dimensions, &scale, &noise_type}; CUresult result = cuLaunchKernel(fill_buffer_function, grid.x, grid.y, grid.z, block.x, block.y, block.z, 0, NULL, params, NULL); checkCudaErrors(cudaDeviceSynchronize()); if (result != CUDA_SUCCESS) { @@ -571,25 +567,6 @@ bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) { 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)); -#ifdef LOG_LEVEL_LOG - - for (int x = 0; x < dimensions.x; ++x) { - for (int y = 0; y < dimensions.y; ++y) { - for (int z = 0; z < dimensions.z; ++z) { - - const unsigned int idx = x + dimensions.x * (y + dimensions.y * z); - - log( " value: " + std::to_string(volume_data_host[idx]), LOG); - - } - } - } - -#endif // LOG_LEVEL_LOG - - - - // create 3D array cudaArray* d_volumeArray = 0; diff --git a/source/gpu_vdb/gpu_vdb.h b/source/gpu_vdb/gpu_vdb.h index d9ce45c..927d675 100644 --- a/source/gpu_vdb/gpu_vdb.h +++ b/source/gpu_vdb/gpu_vdb.h @@ -163,7 +163,7 @@ class GPU_PROC_VOL : public GPU_VDB { __host__ GPU_PROC_VOL(const GPU_PROC_VOL& copy); __host__ ~GPU_PROC_VOL(); - __host__ bool create_volume(float3 min, float3 max, float res); + __host__ bool create_volume(float3 min, float3 max, float res, int noise_type, float scale); private: diff --git a/source/main.cpp b/source/main.cpp index 2e42781..fc8c8a2 100644 --- a/source/main.cpp +++ b/source/main.cpp @@ -1260,7 +1260,9 @@ int main(const int argc, const char* argv[]) // Test procedural volume GPU_PROC_VOL proc_vol; - if(!proc_vol.create_volume(make_float3(0, 100, 0), make_float3(100, 120, 100), 0.25f)) return 0; + float3 proc_box_min = make_float3(0, 1000, 0); + float3 proc_box_max = make_float3(500, 1020, 500); + if(!proc_vol.create_volume(proc_box_min, proc_box_max, 1.0f, 0, 0.1f)) return 0; instances.clear(); instances.push_back(proc_vol); @@ -1458,6 +1460,13 @@ int main(const int argc, const char* argv[]) float exposure = 1.0f; float emission_scale = 0.0f; float emission_pivot = 1.0f; + + + // Noise parameters + int noise_type = 0 , temp_noise_type; + float scale = 1.0f , temp_scale; + float noise_res = 1.0f , temp_res; + // End ImGui parameters @@ -1509,7 +1518,7 @@ int main(const int argc, const char* argv[]) log("Setting up geometry and device pointers...", LOG); float3 center = make_float3(400, 320, -200); - float radius = 100; + float radius = 1; sphere ref_sphere(center, radius); ref_sphere.roughness = 1.0f; ref_sphere.color = make_float3(1.0f, 0 , 0); @@ -1586,6 +1595,13 @@ int main(const int argc, const char* argv[]) earth_atmosphere.m_use_ozone = use_ozone; earth_atmosphere.m_do_white_balance = do_white_balance; earth_atmosphere.m_exposure = exposure; + + // Update temp holders + temp_noise_type = noise_type; + temp_scale = scale; + temp_res = noise_res; + + // Draw imgui //------------------------------------------------------------------- @@ -1620,6 +1636,12 @@ int main(const int argc, const char* argv[]) ImGui::SliderFloat("Camera Aperture", &aperture, .0f, 10.0f); ImGui::Checkbox("Visualize DOF", &viz_dof); + // Noise parameters gui + ImGui::Begin("Camera Parameters"); + ImGui::InputInt("Noise Type", &noise_type, 1); + ImGui::InputFloat("Noise Scale", &scale); + ImGui::InputFloat("Noise resolution", &noise_res); + // Atmosphere Parameters GUI ImGui::Begin("Atmosphere Parameters"); ImGui::SliderFloat("Sky Exposure", &exposure, -10.0f, 10.0f); @@ -1672,6 +1694,19 @@ int main(const int argc, const char* argv[]) } + if (temp_noise_type != noise_type || + temp_res != noise_res || + temp_scale != scale) + { + + proc_vol.create_volume(proc_box_min, proc_box_max, noise_res, noise_type, scale); + check_success(cuMemcpyHtoD(d_volume_ptr, &proc_vol, sizeof(GPU_VDB)* instances.size()) == cudaSuccess); + + kernel_params.iteration = 0; + + } + + // Restart rendering if there is a change if (ctx->change || max_interaction != kernel_params.max_interactions || diff --git a/source/texture_kernels.cu b/source/texture_kernels.cu index 98c9aad..28fd5cb 100644 --- a/source/texture_kernels.cu +++ b/source/texture_kernels.cu @@ -56,6 +56,11 @@ #define INV_PI 1.0f / M_PI +typedef curandStatePhilox4_32_10_t Rand_state; +#define rand(state) curand_uniform(state) + + + 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; @@ -68,7 +73,7 @@ extern "C" __global__ void glow(const Kernel_params kernel_params, float treshol } -extern "C" __global__ void fill_volume_buffer( float *buffer, const int3 dims, const int noise_type) { +extern "C" __global__ void fill_volume_buffer( float *buffer, const int3 dims, const float scale, const int noise_type) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -78,5 +83,46 @@ extern "C" __global__ void fill_volume_buffer( float *buffer, const int3 dims, c const unsigned int idx = x + dims.x * (y + dims.y * z); - buffer[idx] = 1.0f; + Rand_state rand_state; + + int seed = 123; + float du = 1.0f / (float)dims.x; + + float dx = cudaNoise::randomFloat(482 + floor(rand(&rand_state) * 2) * 47 + seed) / (float)dims.x; + float dy = cudaNoise::randomFloat(472 + floor(rand(&rand_state) * 2) * 38 + seed) / (float)dims.y; + float dz = cudaNoise::randomFloat(348 + floor(rand(&rand_state) * 2) * 14 + seed) / (float)dims.z; + + float3 pos = make_float3(x+dx, y+dy, z+dz); + + switch (noise_type) + { + case(0): + buffer[idx] = cudaNoise::perlinNoise(pos, scale, seed); + break; + case(1): + buffer[idx] = cudaNoise::simplexNoise(pos, scale, seed); + break; + case(2): + buffer[idx] = cudaNoise::worleyNoise(pos, scale, seed, 300.1f, 4, 4, 1.0f); + break; + case(3): + buffer[idx] = cudaNoise::repeaterPerlin(pos, scale, seed, 128, 1.9f, 0.5f); + break; + case(4): + buffer[idx] = cudaNoise::repeaterPerlinAbs(pos, scale, seed, 128, 1.9f, 0.5f); + break; + case(5): + buffer[idx] = cudaNoise::fractalSimplex(pos, scale, seed, du, 512, 1.5f, 0.95f); + break; + case(6): + buffer[idx] = cudaNoise::repeaterTurbulence(pos, 0.2f, scale, seed, 0.8f, 32, cudaNoise::BASIS_PERLIN, cudaNoise::BASIS_PERLIN); + break; + case(7): + buffer[idx] = cudaNoise::cubicValue(pos, scale, seed); + break; + case(8): + buffer[idx] = cudaNoise::spots(pos, scale, seed, 0.1f, 0, 8, 1.0f, cudaNoise::SHAPE_STEP); + break; + } + } \ No newline at end of file