will allocate and fill procedural volume data in texture_kernels.cu

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;
+}

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_

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);

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;
+
 }