Skip to content

Commit

Permalink
pass the kernel functions in class function. Test volume with a new gui
Browse files Browse the repository at this point in the history
  • Loading branch information
@sergeneren
sergeneren committed Dec 27, 2019
1 parent 1251533 commit c3f853d
Show file tree
Hide file tree
Showing 4 changed files with 88 additions and 30 deletions.
27 changes: 2 additions & 25 deletions source/gpu_vdb/gpu_vdb.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down Expand Up @@ -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);
Expand All @@ -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) {
Expand All @@ -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;

Expand Down
2 changes: 1 addition & 1 deletion source/gpu_vdb/gpu_vdb.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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:

Expand Down
39 changes: 37 additions & 2 deletions source/main.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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


Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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
//-------------------------------------------------------------------

Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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 ||
Expand Down
50 changes: 48 additions & 2 deletions source/texture_kernels.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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;
Expand All @@ -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;
}

}

0 comments on commit c3f853d

Please sign in to comment.