DEM simulation accessing collision data and cutil.h

I was wondering how I go about accessing the collision data from the program. I am trying to come up with a simple input file so I can at least get going on my algorithms a little bit and am not sure how the values are stored and the syntax to access them.

Also, I have been unable to compile the project provided online because it keeps telling me cutil.h is missing. I have installed everything from NVIDIA's website as well as followed the directions for compiling code in the slides, but it just won't compile. What do I need to do?

Thanks,

Ben

I should also add that I would like to know the syntax for accessing each element in the float4. Furthermore, is it possible to set all the x components equal to a pointer, say *px, so that it might be easier to use the cudaMemcpy to transfer the data into the GPU?

Thanks.

See the post here (viewtopic.php?f=11&t=152#p834) for some info about the collision data. The array contactDataD is already on the device after collision detection.

See the thread here (viewtopic.php?f=11&t=167) for info about float4. If the variable temp is of type float4, then you access data like temp.x, temp.y, temp.z, temp.w.

See the post here (viewtopic.php?f=11&t=152#p903) for info about the build issue. Respond in that thread if you still cannot get it to work.



So, as an example you can have the following:

Kernel:

  Code:

__global__ void someKernel(float4* PosD, contact_Data* CData)
{
int thid = threadIdx.x; //each thread processes data related to contact number thid
int iA=contactDataD[thid].objectIdA; //the ID number of the first body in the contact
int iB=contactDataD[thid].objectIdB; //the ID number of the second body in the contact
float4 posA = PosD[iA]; //the position and radius data for body A
float4 posB = PosD[iB]; //the position and radius data for body B

float posAx = posA.x; //the x position of body A
float posBx = posB.x; //the x position of body B

float radiusA = posA.w; //the radius of body A
float radiusB = posB.w; //the radius of body B
}

CPU code:

  Code:

void ParticleSystem::computeForces(float delta_t)
{
//define nB, nT appropriately
someKernel<<<nB, nT>>>(bodyPosD, contactDataD); //note that position data and contact data are already on the device
}


I don't understand how you passed contactDataD to the function. What is CData? What is contactData*?

Should it be CData[tid].objectIdA? It seems like you are declaring a variable for the kernel and then never using it, or at least using the incorrect variable when declaring iA and iB.

You are right, I had a mistake in the kernel code. It should be like this:

  Code:

__global__ void someKernel(float4* PosD, contact_Data* CData)
{
int thid = threadIdx.x; //each thread processes data related to contact number thid
int iA=CData[thid].objectIdA; //the ID number of the first body in the contact
int iB=CData[thid].objectIdB; //the ID number of the second body in the contact
float4 posA = PosD[iA]; //the position and radius data for body A
float4 posB = PosD[iB]; //the position and radius data for body B

float posAx = posA.x; //the x position of body A
float posBx = posB.x; //the x position of body B

float radiusA = posA.w; //the radius of body A
float radiusB = posB.w; //the radius of body B
}

Remember that contactDataD is a pointer to variable of type contact_Data. The kernel should take as an argument a pointer to a variable of type contact_Data. The kernel should work when passing any pointer to collision data. The definition of the kernel, ( __global__ void someKernel(float4* PosD, contact_Data* CData) ) shows that the kernel takes two arguments, where the first one is a pointer to float4 and the second is a pointer to contact_Data. Here, contact_Data* means "pointer to variable of type contact_Data" and CData means "this pointer is called by the name CData in the kernel".

Where is the mass of each particle stored? I was going to calculate the effective mass and cannot find it. Also, how do I access the force constants (k and gamma) in order to calculate force? Lastly, is there a source that explains how to use the sortbykey?

Thanks,

Ben

Look in particleSystem.h. All the variables you need to use are declared there (I think). The spring and damping coefficients are there (float kn, float gn). Also, the particle mass (float pMass). Note that this is a single constant which can be used for all bodies. In this case, the effective mass is m_eff=pMass/2. If you want, you could generalize it to a vector so the particles can have unique masses.
These three variables (kn, gn, pMass) are initialized on the CPU and not pre-copied to the GPU. The easiest way to use them is to pass them as arguments to the appropriate kernel:

  Code:

__global__ void someKernel(float4* PosD, contact_Data* CData, float m, float k, float g)
{
int thid = threadIdx.x; //each thread processes data related to contact number thid
int iA=contactDataD[thid].objectIdA; //the ID number of the first body in the contact
int iB=contactDataD[thid].objectIdB; //the ID number of the second body in the contact
//...
}

void ParticleSystem::computeForces(float delta_t)
{
//define nB, nT appropriately
someKernel<<<nB, nT>>>(bodyPosD, contactDataD, pMass, kn, gn); //note that position data and contact data are already on the device
}


sort_by_key is a Thrust function. You can look at the documentation online (http://code.google.com/p/thrust/wiki/Qu ... de#Sorting).
The idea is that you call sort_by_key with three arguments. The first points to the first key, the second points one past the last key, and the third points to the first value.
Here is a sample code:

  Code:

#include <thrust/sort.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/constant_iterator.h>
#include <cuda.h>
#include <cutil.h>


int main(int argc, char* argv[])
{
   //The data on the Host
   const int N = 8;
   int keyH[N] = {3, 1, 5, 3, 7, 3, 2, 9};
   float valueH[N] = {4.3, 9.0, 13.2, -4.2, 3.3, 20.1, -4.3, 4.1};

   //Show the data:
   printf("Data Before Sorting:\n");
   for(int i=0; i<N; i++)
   {
      printf("%i->%f\n",keyH[i], valueH[i]);
   }

   //Allocate device memory and copy data to device
   int * key_ptr;
   cudaMalloc((void **) &key_ptr, N*sizeof(int));
   cudaMemcpy(key_ptr, keyH, N*sizeof(int), cudaMemcpyHostToDevice);

   float * value_ptr;
   cudaMalloc((void **) &value_ptr, N*sizeof(float));
   cudaMemcpy(value_ptr, valueH, N*sizeof(float), cudaMemcpyHostToDevice);
   
   //This is similar to the place you would be in your code, with some data already on the device
   // wrap device memory pointers with thrust device_ptr to use in thrust functions
   thrust::device_ptr<int> key_ptr_thrust(key_ptr);
   thrust::device_ptr<float> value_ptr_thrust(value_ptr);

   //Call the sort function
   thrust::sort_by_key(key_ptr_thrust, key_ptr_thrust+N, value_ptr_thrust);

   //Copy result back to CPU
   cudaMemcpy(keyH, key_ptr, N*sizeof(int), cudaMemcpyDeviceToHost);
   cudaMemcpy(valueH, value_ptr, N*sizeof(float), cudaMemcpyDeviceToHost);

   //Show the data, now sorted:
   printf("Data After Sorting:\n");
   for(int i=0; i<N; i++)
   {
      printf("%i->%f\n",keyH[i], valueH[i]);
   }

   cudaFree(key_ptr);
   cudaFree(value_ptr);
   
   return 0;
}


I haven't yet looked at reduce_by_key yet, but I will do so tomorrow morning and hopefully put up a complete demo code with both functions soon.

This seems pretty straight forward, but one thing I am having trouble understanding is when you say "wrap device memory pointers with thrust device_ptr to use in thrust functions. I was wondering what is going on with this segment of code:

thrust::device_ptr<int> key_ptr_thrust(key_ptr);
thrust::device_ptr<float> value_ptr_thrust(value_ptr);

Then, after you call sort_by_key, are the pointers that are on the device sorted?

We have the normal device pointers (key_ptr, value_ptr). These are just like you always use for GPU global memory.
Thrust uses its own type of pointers, which can have some more hidden information. For our purposes, we can make a thrust pointer from our device pointer.

thrust::device_ptr<int> key_ptr_thrust(key_ptr);
-This means create a Thrust pointer to device memory which holds ints (note that int matches the type used when creating key_ptr).
-The thrust pointer is called key_ptr_thrust
-It is created from the device pointer key_ptr

The call to the sort function is made with the thrust pointers, but 'under the hood' it actually uses the device pointer we used to create the thrust pointer. Therefore, the device pointer still points to the same memory location (and the array is now sorted).

Here is some code that I _think_ should work for the reduce_by_key. I am having some issues with my GPU so I can't test it, but I hope it works. You can add this code to the code I posted before after the sorting step.
I will try to test it tomorrow and correct it if it isn't right.

Other Thrust resources:
http://code.google.com/p/thrust/wiki/QuickStartGuide
http://forums.nvidia.com/index.php?showtopic=163877

  Code:

   //Allocate some more memory to hold the result of the reduce
   int * keyOut;
   cudaMalloc((void **) &keyOut, N*sizeof(int));
   thrust::device_ptr<int> keyOut_thrust(keyOut);
   
   float * valueOut;
   cudaMalloc((void **) &valueOut, N*sizeof(float));
   thrust::device_ptr<float> valueOut_thrust(valueOut);

   //Call the reduce function
   //int N_afterReduce=0;
   thrust::pair<thrust::device_ptr<int>,thrust::device_ptr<float> > new_end;
   new_end = thrust::reduce_by_key(key_ptr_thrust, key_ptr_thrust+N, value_ptr_thrust, keyOut_thrust, valueOut_thrust);
   int N_afterReduce = (new_end.first.get()-keyOut_thrust.get());

   printf("\nNumber of unique keys: %i\n\n",N_afterReduce);

   //Copy result back to CPU
   cudaMemcpy(keyH, keyOut, N_afterReduce*sizeof(int), cudaMemcpyDeviceToHost);
   cudaMemcpy(valueH, valueOut, N_afterReduce*sizeof(float), cudaMemcpyDeviceToHost);

   //Show the data, now sorted and reduced:
   printf("\nData After Sorting and Reduce:\n");
   for(int i=0; i<N_afterReduce; i++)
   {
      printf("%i->%f\n",keyH[i], valueH[i]);
   }