Since NVIDIA has this topic I think its not bad putting here my 4:30 hours long lecture at a Summer School. Please leave your comments inside here. I will be tracking this topic.
http://www.youtube.com/watch?v=ykUs4MYOwcY&feature=youtu.be
This video is for educational purposes and should not be used for profit.
Just a nice image that what I teach actually work…this is the adaptive octree on the Stanford bunny. Its a 5 level octree and levels 1, 3, 5 are displayed. I still have the privilege to develop what I think…
External Media
Just a note. Reading from the specific paper having only as friend Google search is a torture in the brain. Only people with strong nerves should attempt it and moreover code it…A little gossip. :-)
I am also tempted to include the host code that builds this nice image. The kernels to do it are left as an exercise. I learned from this Forum a nice way to use double pointers. I believe we should invest a lot in demonstrating not to code lame code. This is why I am actually devoting some time to contribute in this Forum.
dev_octree_node** dnodeArrayd;
cudaMalloc((void**)& dnodeArrayd, (D+1)*sizeof(dev_octree_node*));
dev_octree_node** supportNode = (dev_octree_node**) malloc((D+1)*sizeof(dev_octree_node*));
cudaMemcpy(supportNode, dnodeArrayd, (D+1)*sizeof(dev_octree_node*), cudaMemcpyDeviceToHost);
dev_octree_node* dnodeArray;
int* dnodeBaseAddress;
int* startDepthIndex = (int*) malloc((D+1)*sizeof(int));
int totalNumbNodes = 0;
for (int d = D; d >= 0; d--){
cudaMalloc((void**)& dnodeBaseAddress, numbUniqueNodes*sizeof(int));
get_numbers_base_address(duniqueNodes, dnodeBaseAddress, numbUniqueNodes);
thrust::device_ptr<int> thrstdbaseAddress = thrust::device_pointer_cast(dnodeBaseAddress);
thrust::inclusive_scan(thrstdbaseAddress, thrstdbaseAddress + numbUniqueNodes, thrstdbaseAddress);
int numbNodesd = d > 0 ? thrstdbaseAddress[numbUniqueNodes - 1] + 8 : 1;
cudaMalloc((void**)& supportNode[D-d], numbNodesd*sizeof(dev_octree_node));
populate_node_array_from_base_address(duniqueNodes, supportNode[D-d], dnodeBaseAddress, numbUniqueNodes, numbNodesd, currentNodeDepth);
cudaFree(duniqueNodes);
cudaFree(dnodeBaseAddress);
numbUniqueNodes = numbNodesd / 8;
if (d > 0){
cudaMalloc((void**)&duniqueNodes, numbUniqueNodes*sizeof(dev_octree_node));
construct_next_level(supportNode[D-d], duniqueNodes, numbNodesd);
}
startDepthIndex[D-d] = totalNumbNodes;
totalNumbNodes += numbNodesd;
currentNodeDepth *= 2.0;
}
cudaMalloc((void**)& dnodeArray, totalNumbNodes*sizeof(dev_octree_node));
for (int i = 0; i <= D; i++){
if (i < D) cudaMemcpy(dnodeArray+startDepthIndex[i], supportNode[i], (startDepthIndex[i+1]-startDepthIndex[i])*sizeof(dev_octree_node), cudaMemcpyDeviceToDevice);
else
cudaMemcpy(dnodeArray+startDepthIndex[i], supportNode[i], sizeof(dev_octree_node), cudaMemcpyDeviceToDevice);
cudaFree(supportNode[i]);
}
cudaFree(dnodeArrayd);
free(supportNode);
The code to build the Look Up Table for the neighbor information is this :
int c[6][6][6];
int p[6][6][6];
int LUTparent[8][27];
int LUTchild[8][27];
int numbParent = 0;
for (int k = 5; k >= 0; k -= 2)
for (int i = 0; i < 6; i += 2)
for (int j = 5; j >= 0; j -= 2){
int numb = 0;
for (int l = 0; l < 2; l++)
for (int m = 0; m < 2; m++)
for (int n = 0; n < 2; n++){
c[i+m][j-n][k-l] = numb++;
p[i+m][j-n][k-l] = numbParent;
}
numbParent++;
}
int numbLUT = 0;
for (int k = 3; k > 1; k--)
for (int i = 2; i < 4; i++)
for (int j = 3; j > 1; j--){
int numb = 0;
for (int n = 1; n >= -1; n--)
for (int l = -1; l <= 1; l++)
for (int m = 1; m >= -1; m--){
LUTparent[numbLUT][numb] = p[i+l][j+m][k+n];
LUTchild[numbLUT][numb++] = c[i+l][j+m][k+n];
}
numbLUT++;
}
You can now build with my Lecture a very efficient algorithm for Marching Cubes. It can be the theory of H’oppe’s (I am putting the necessary tone :-) approach, you have the right cubes. Good Luck.
