Parallel Poisson Disk Sampling on Surfaces

Become a member of the CGSociety

Connect, Share, and Learn with our Large Growing CG Art Community. It's Free!

THREAD CLOSED
 
Thread Tools Search this Thread Display Modes
Old 09 September 2013   #1
Parallel Poisson Disk Sampling on Surfaces

Hi,

I present my implementation of algorithm explained in "Parallel Poisson Disk Sampling with Spectrum Analysis on Surfaces" document. Program uniformly scatters points on a mesh.
Full documentation available here: http://research.microsoft.com/apps/....aspx?id=135760
Project was written as NVIDIA CUDA based plugin (DG Node) to Autodesk Maya 2013.





Scattering points is almost real-time. In comparison to sequential CPU algorithms speedup is significant. This parallel scatter works well (fast) with high polycount meshes.

Plugin compiled for 64bit Linux, Autodesk Maya 2013, Funtoo/Gentoo distribution, CUDA libraries required. Haven't tested it on other distributions. If I find time I'll try to compile it for Windows. Calculation results can be connected to instancer node. No rotations computed, so practical use-case is limited at the moment, but I'm going to develop more complex tool soon.

Greetings!
 
Old 09 September 2013   #2
Very interesting! Does it work with nonmanifold meshes (T-shapes eg.)? And if you have time, a windows OS port of this plugin would be nice!
 
Old 09 September 2013   #3
Originally Posted by zaskar: Very interesting! Does it work with nonmanifold meshes (T-shapes eg.)? And if you have time, a windows OS port of this plugin would be nice!


DG Node operates on mesh triangles, so it should work with non-manifold meshes. I implemented sampling with euclidean distance, so if there are two triangles described on same three points, final result will be like only one triangle receive samples. Absolute distance between two samples cannot be smaller than r.
I'll try to compile it for Win.
 
Old 09 September 2013   #4
Intersting! but what would you that for? is it particles to which you can instance mesh to it?
Also would you do a compile for maya 2014.

cheers


edit: sorry you already said you could link instancer to it. would be nice to be able to get some more data out of it. Uv, position,normal of object.
 
Old 09 September 2013   #5
Originally Posted by thematt: Intersting! but what would you that for? is it particles to which you can instance mesh to it?
Also would you do a compile for maya 2014.

cheers


edit: sorry you already said you could link instancer to it. would be nice to be able to get some more data out of it. Uv, position,normal of object.


Currently it outputs array of points. This DG node can be used for example to uniformly scatter grass on ground, another instance is an input to remesher, etc. Good base to do something more, and it's good it is really fast. When I started to write this plugin I intended to create feather tool and I hope it'll happen in future.
It is possible to outputs mesh normals, barycentric coordinates on tri & tri ID, UV.
 
Old 09 September 2013   #6
Wow that's really cool. Would love a 2013 or 2014 Win64 version compiled - could use it right now.

Currently I'm using a particle emitter with emission overlap pruning, but it takes ages (hours) to get a good even distribution.
 
Old 09 September 2013   #7
This is looking great! Especially the speed of the whole things is amazing.

You mentioned that outputing the normals is possible. Are you doing that from CUDA as well or are you getting the normals with the Maya API. I've had the API becoming extremely slow when getting a large amount of normals (and binormals and tangents) at arbitrary locations on the mesh.

And yeah, I should really look into how I can make plug-ins for Maya that run on CUDA. If you ever feel like making a basic tutorial I would definitely be interested. Hehe.

-Roy
 
Old 09 September 2013   #8
Wow that's really cool. Would love a 2013 or 2014 Win64 version compiled - could use it right now.
Currently I'm using a particle emitter with emission overlap pruning, but it takes ages (hours) to get a good even distribution.

This is looking great! Especially the speed of the whole things is amazing.

Thanks! I'll try to compile it for Win later.

You mentioned that outputting the normals is possible. Are you doing that from CUDA as well or are you getting the normals with the Maya API. I've had the API becoming extremely slow when getting a large amount of normals (and binormals and tangents) at arbitrary locations on the mesh.

Computing normals can be performed on GPU too, Maya API is used to get mesh data: vertices positions, vertices ID in triangle. Computing normals for triangles is easy, because calculations are independent, but if you want to get normal at point location that is interpolated among surrounding faces case is more complicated but it's still possible I think.

And yeah, I should really look into how I can make plug-ins for Maya that run on CUDA. If you ever feel like making a basic tutorial I would definitely be interested. Hehe.

OK here is short advice how to compile CUDA plug-in in Linux environment.

My template for DG Node that utilizes CUDA:


     
     /*
      * mayaCuda.cu
      *
      *  Created on: May 28, 2013
      *	  Author: michal
      */
     
     #ifdef __CDT_PARSER__
     #define __global__
     #define __device__
     #define __shared__
     #endif
     
     #include <cuda_runtime.h>
     
     __global__ void
     vectorAdd(const float *A, float *C, int numElements)
     {
     	int i = blockDim.x * blockIdx.x + threadIdx.x;
     
     	if (i < numElements)
     	{
     		C[i] = A[i] + 1.0f;
     	}
     }
     
     
     #define LINUX
     
     #include <iostream>
     
     #include <maya/MGlobal.h>
     #include <maya/MStatus.h>
     #include <maya/MPxNode.h>
     #include <maya/MFnNumericAttribute.h>
     #include <maya/MFnMatrixAttribute.h>
     #include <maya/MFnPlugin.h>
     #include <maya/MFloatMatrix.h>
     #include <maya/MScriptUtil.h>
     #include <maya/MFnMesh.h>
     
     
     #define MAKE_INPUT(attr)		\
     	CHECK_MSTATUS(attr.setKeyable(true) );			  \
     	CHECK_MSTATUS( attr.setStorable(true) );	\
     	CHECK_MSTATUS( attr.setReadable(true) );	\
     	CHECK_MSTATUS( attr.setWritable(true) );
     
     #define MAKE_OUTPUT(attr)					   \
     	CHECK_MSTATUS( attr.setKeyable(false) );	\
     	CHECK_MSTATUS( attr.setStorable(false) );   \
     	CHECK_MSTATUS( attr.setReadable(true) );	\
     	CHECK_MSTATUS( attr.setWritable(false) );
     
     
     using namespace std;
     
     class mayaCuda : public MPxNode
     {
     	public:
     		mayaCuda();
     		virtual ~mayaCuda();
     
     		static  void* creator();
     		static  MStatus initialize();
     		virtual MStatus compute(const MPlug& plug, MDataBlock& dataBlock);
     
     	public:
     		static  MTypeId	 id;
     
     		static	MObject		input;
     		static	MObject		output;
     
     };
     
     MTypeId	 mayaCuda::id( 0x000A9 );
     
     MObject	 mayaCuda::input;
     MObject	 mayaCuda::output;
     
     mayaCuda::mayaCuda() {}
     mayaCuda::~mayaCuda() {}
     
     void* mayaCuda::creator()
     {
     		return new mayaCuda();
     }
     
     MStatus mayaCuda::initialize()
     {
     		MStatus stat;
     
     		MFnNumericAttribute numAttr;
     
     		input = numAttr.create("input", "in", MFnNumericData::kFloat);
     		MAKE_INPUT(numAttr);
     		numAttr.setDefault(0.0);
     		stat = addAttribute(input);
     		CHECK_MSTATUS(stat);
     
     		output = numAttr.create("output", "out", MFnNumericData::kFloat);
     		MAKE_OUTPUT(numAttr);
     		numAttr.setDefault(0.0);
     		stat = addAttribute(output);
     		CHECK_MSTATUS(stat);
     
     		attributeAffects( mayaCuda::input, mayaCuda::output);
     
     		return MStatus::kSuccess;
     }
     
     MStatus mayaCuda::compute(const MPlug& plug, MDataBlock& block)
     {
     		cudaError_t err = cudaSuccess;
     		MDataHandle handle;
     
     		bool plugFound = false;
     
     		if(plug == output)
     		{
     			cout << endl;
     			plugFound = true;
     
     			float inputF;
     
     			handle = block.inputValue(input);
     			inputF = handle.asFloat();
     
     			int numElements = 1;
     			size_t size = numElements * sizeof(float);
     
     			float *h_A = (float *)malloc(size);
     			float *h_C = (float *)malloc(size);
     
     			h_A[0] = inputF;
     
     			cout << "inputF: " << inputF << endl;
     
     			float *d_A = NULL;
     			err = cudaMalloc((void **)&d_A, size);
     
     			if (err != cudaSuccess)
     			{			  
     				cout << "Failed to allocate device vector A, error code: " << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     
     			float *d_C = NULL;
     			err = cudaMalloc((void **)&d_C, size);
     
     			if (err != cudaSuccess)
     			{				
     				cout << "Failed to allocate device vector C, error code: " << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     		   
     			//cout << "Copy input data from the host memory to the CUDA device" << endl;
     			err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
     
     			if (err != cudaSuccess)
     			{			   
     				cout << "Failed to copy vector A from host to device, error code: " << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     
     			// Launch the Vector Add CUDA Kernel
     			int threadsPerBlock = 32;
     			int blocksPerGrid =1;
     		   
     			//cout << "CUDA kernel launch" << endl;
     
     			vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_C, numElements);
     			err = cudaGetLastError();
     
     			if (err != cudaSuccess)
     			{			   
     				cout << "Failed to launch vectorAdd kernel, error code: " << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     		   
     			//cout << "Copy output data from the CUDA device to the host memory" << endl;
     			err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
     
     			if (err != cudaSuccess)
     			{			   
     				cout << "Failed to copy vector C from device to host, error code: " << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     
     			err = cudaFree(d_A);
     
     			if (err != cudaSuccess)
     			{			   
     				cout << "Failed to free device vector A, error code: " << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     
     			err = cudaFree(d_C);
     
     			if (err != cudaSuccess)
     			{			   
     				cout << "Failed to free device vector C (error code %s)!\n" << cudaGetErrorString(err);
     				return MStatus::kFailure;
     			}
     
     
     			float result = h_C[0];
     
     			cout << "result: " << result << endl;
     
     			handle = block.outputValue(output);
     			handle.set(result);
     
     
     			// Free host memory
     			free(h_A);
     			free(h_C);
     
     			// Reset the device and exit
     			err = cudaDeviceReset();
     
     			if (err != cudaSuccess)
     			{
     				cout << "Failed to deinitialize the device! error=" << cudaGetErrorString(err) << endl;
     				return MStatus::kFailure;
     			}
     
     		}
     
     		if(plugFound)
     		{
     			block.setClean(plug);
     			return MStatus::kSuccess;
     		}
     
     		return MStatus::kUnknownParameter;
     
     }
     
     
     MStatus initializePlugin( MObject obj )
     {
     		MStatus result;
     		MFnPlugin plugin( obj, PLUGIN_COMPANY, "3.0", "Any");
     		result = plugin.registerNode( "mayaCuda", mayaCuda::id, mayaCuda::creator, mayaCuda::initialize, MPxNode::kDependNode );
     		cout << "mayaCuda initialized" << endl;
     		return result;
     }
     
     MStatus uninitializePlugin( MObject obj)
     {
     		MStatus result;
     		MFnPlugin plugin( obj );
     
     		result = plugin.deregisterNode( mayaCuda::id );
     		return result;
     }
     
     

Code above simply adds on GPU 1.0 to input.

First thing, to compile Maya plugin that uses CUDA you need to modify header files and rename some data types, because both APIs define something like int3, float3 structures.
Next, there is conflict between GCC compilers. Autodesk Maya API requires GCC 4.1.2, but to compile CUDA newer version of GCC is needed. I handle that by splitting compilation into few steps:


     
     nvcc -cuda -arch compute_11 mayaCuda.cu --compiler-bindir=/usr/x86_64-pc-linux-gnu/gcc-bin/4.1.2/ -I/home/michal/dev/include/ -I/opt/cuda/sdk/common/inc/ -D_BOOL -DREQUIRE_IOSTREAM
     
     g++-4.1.2 mayaCuda.cu.cpp.ii -O3 -Wall -c -m64 -pthread -D_BOOL -DLINUX_64 -DREQUIRE_IOSTREAM -fPIC -mtune=nocona
     
     g++-4.1.2 -l:libstdc++.so.6.0.16 -L/usr/lib64/gcc/x86_64-pc-linux-gnu/4.6.3 -L/opt/cuda/lib64 -lcudart -lcuda -L/usr/autodesk/maya2013-x64/lib -shared -m64 -O3 -pthread -pipe -D_BOOL -DLINUX -DREQUIRE_IOSTREAM -mtune=nocona -Wl,-Bsymbolic -lOpenMaya -lFoundation -lcurand -o "mayaCuda.so"  ./mayaCuda.cu.cpp.o
     
     cp mayaCuda.so ~/maya/2013-x64/plug-ins/mayaCuda.so
     
     

I know that this description isn't too clear, but I hope it'll help in some way.
 
Old 10 October 2013   #9
Hello, any chances to see a windows compile in the near future?
 
Old 10 October 2013   #10
Windows build for Maya 2013 with this node and other tools should be released in two weeks - I hope so. I'll leave you message in this thread.
 
Old 10 October 2013   #11
that s really cool stuff man , any chance to have some banch marks?

like time a couple of meshes at different samplings and ofc spec of the machine?
__________________
My website/blog , with a lot of tips&tricks about maya API,python,and pyqt , and tutorials
www.marcogiordanotd.com
llinkedin
 
Old 11 November 2013   #12
Originally Posted by zaskar: Hello, any chances to see a windows compile in the near future?


It was predictable: it took longer than I expected. Additional content (scripts, tutorial, graphics) was really time-consuming. Please find Windows build in this thread: http://forums.cgsociety.org/showthr...?f=89&t=1139186

Originally Posted by giordi: that s really cool stuff man , any chance to have some banch marks? like time a couple of meshes at different samplings and ofc spec of the machine?


Apologize for late response. I was totally absorbed by my project. Examples presented here were created using GEFORCE GTX 660M with 2GB of memory and 384 CUDA cores. Please find Windows build at my website. You can check performance by yourself :-)
 
Old 11 November 2013   #13
Thread automatically closed

This thread has been automatically closed as it remained inactive for 12 months. If you wish to continue the discussion, please create a new thread in the appropriate forum.
 
Thread Closed share thread



Posting Rules
You may not post new threads
You may not post replies
You may not post attachments
You may not edit your posts

vB code is On
Smilies are On
[IMG] code is On
HTML code is Off
CGSociety
Society of Digital Artists
www.cgsociety.org

Powered by vBulletin
Copyright 2000 - 2006,
Jelsoft Enterprises Ltd.
Minimize Ads
Forum Jump
Miscellaneous

All times are GMT. The time now is 10:49 AM.


Powered by vBulletin
Copyright ©2000 - 2017, Jelsoft Enterprises Ltd.