testing procedural volume

source/gpu_vdb/gpu_vdb.cpp

@@ -507,13 +507,21 @@ 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) {
 
+	log("Creating procedural volume...", LOG);
+
 	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));
+
+#ifdef LOG_LEVEL_LOG
+	log("XForm: ", LOG);
+	xform.print();
+#endif // LOG_LEVEL_LOG
+
 	set_xform(xform);
 
 	int dim_x = floorf((max.x - min.x) / res);
@@ -524,8 +532,8 @@ bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) {
 
 	// 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.bmin = min;
+	vdb_info.bmax = make_float3(min.x + dim_x, min.y + dim_y, min.z + dim_z);
 	vdb_info.voxelsize = res;
 	vdb_info.min_density = .0f;
 	vdb_info.max_density = 1.0f;
@@ -535,25 +543,53 @@ bool GPU_PROC_VOL::create_volume(float3 min, float3 max, float res) {
 	// 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);
 	checkCudaErrors(cudaMalloc(&device_density_buffer, dimensions.x * dimensions.y * dimensions.z * sizeof(float)));
-
-	dim3 block(16, 16, 16);
+	
+	dim3 block(8, 8, 8);
 	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);
-
+	log("filling volume buffer in device...", LOG);
+	void* params[] = { &device_density_buffer , &dimensions, &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) {
+		printf("Unable to launch fill_buffer_function! result: %i\n", result);
+		return false;
+	}
+
 	// send buffer to texture 
 
 	cudaExtent vol_size;
 	vol_size.width = dim_x;
 	vol_size.height = dim_y;
 	vol_size.depth = dim_z;
 
+	log("transport volume buffer from device to host...", LOG);
 	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;
 

source/main.cpp

@@ -1258,13 +1258,19 @@ int main(const int argc, const char* argv[])
 	error = cuModuleGetFunction(&cuTestGeometryKernel, Module, "test_geometry_list");
 	if (error != CUDA_SUCCESS) log("cuModuleGetFunction " + error, ERROR);
 
+	// 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;
 
+	instances.clear();
+	instances.push_back(proc_vol);
+
 	// Send volume instances to gpu
 
 	CUdeviceptr d_volume_ptr;
 	check_success(cuMemAlloc(&d_volume_ptr, sizeof(GPU_VDB) * instances.size()) == cudaSuccess);
 	check_success(cuMemcpyHtoD(d_volume_ptr, instances.data(), sizeof(GPU_VDB) * instances.size()) == cudaSuccess);
-
+	
 
 	// Create BVH from vdb vector 
 	log("Creating BVH and Octree structures...", LOG);

source/texture_kernels.cu

@@ -68,7 +68,7 @@ extern "C" __global__ void glow(const Kernel_params kernel_params, float treshol
 
 }
 
-extern "C" __global__ void fill_volume_buffer(float *buffer, int3 dims, int noise_type) {
+extern "C" __global__ void fill_volume_buffer( float *buffer, const int3 dims, const int noise_type) {
 
 	int x = blockIdx.x * blockDim.x + threadIdx.x;
 	int y = blockIdx.y * blockDim.y + threadIdx.y;
@@ -78,6 +78,5 @@ extern "C" __global__ void fill_volume_buffer(float *buffer, int3 dims, int nois
 
 	const unsigned int idx = x + dims.x * (y + dims.y * z);
 
-	buffer[idx] = .0f;
-
+	buffer[idx] = 1.0f;
 }