Ok, so previously I was using atomicAdd and atomicExch in my CUDA code and it wouldn’t compile. In order to get my code to compile I took the following step
However, since then I’ve been getting weird results from my code, particularly the section using the atomic functions. So I added the following line into
my code
#ifndef CUDA_NO_SM_11_ATOMIC_INTRINSICS
std::cout << "Wait, I'm not using atomics?" << std::endl;
#endif
Since I saw the above text print out on my console, does this mean that I’m not actually using atomic operations?
I’m using atomics in an SPH simulation and during grid generation I need to add the particle id, and increment the counter (this is one of the methods described in the CUDA SDK). In the manner shown below.
So clearly if atomics aren’t working then I would see some pretty bizarre results when i print out the contents of grid cells (which is what I’ve been seeing).
Anyway, I’m just hoping someone might have some insight into the problem. Thanks for looking. :)
That atomicExch line looks fishy. What are you hoping to accomplish with it? The fact that you don’t use the return value (the value you exchanged out) means that you don’t need an atomic there at all.
A more common use of atomics for assembling data is to do an atomic add (or atomic inc) first, then use the returned incremented values as an index to safely write an ID or data into a global array.
Without a little more context, it’s hard to see what you’re trying to do, and therefore hard to help.
I apologize for the vague bit of code I posted. In this section of my code I am inserting a particle into a 3d grid by its location (x, y, z). So once I have determined which cell of my grid a particle belongs to, I insert that particle’s id into that grid cell and increment the counter for that particular grid cell. I used atomicExch only because I didn’t see any other way to replace the value & guarantee that no other block would be doing the same.
Here is an abbreviated version of my update grid kernel:
/=====================================================================================================
// UPDATE GRID
//=====================================================================================================
__global__ void Update(particle * dev_particles, int * dev_grid, int * dev_counters)
{
int tid = blockIdx.x;
int x, y, z;
if(tid < sph_params.np)
{
x = dev_particles[tid].position.x;
y = dev_particles[tid].position.y;
z = dev_particles[tid].position.z;
//bounds check
if(x >= sph_params.grid_x_range)
{
x = sph_params.grid_x_range - 1;
}
if(x < 0 )
{
x = 0;
}
if(y >= sph_params.grid_y_range)
{
y = sph_params.grid_y_range - 1;
}
if(y < 0 )
{
y = 0;
}
if(z >= sph_params.grid_z_range)
{
z = sph_params.grid_z_range - 1;
}
if(z < 0 )
{
z = 0;
}
//if we have not filled up that particular grid cell - add the particle id and increment that cell's counter
if(dev_counters[x * sph_params.grid_y_range * sph_params.grid_z_range + y * sph_params.grid_z_range + z] < sph_params.max_cell)
{
atomicExch(&dev_grid[ x * sph_params.grid_y_range * sph_params.grid_z_range * sph_params.max_cell + y * sph_params.grid_z_range * sph_params.max_cell + z * sph_params.max_cell + dev_counters[x * sph_params.grid_y_range * sph_params.grid_z_range + y * sph_params.grid_z_range + z] ], dev_particles[tid].id);
atomicAdd( &dev_counters[x * sph_params.grid_y_range * sph_params.grid_z_range + y * sph_params.grid_z_range + z], 1);
}
}
}
I am attempting to implement the method described in the CUDA SDK Particle example .pdf, and i quote
UPDATE: Ok, I see my problem now, I need to be launching a thread for every particle… doh
I just wanted to post what I did to fix the problem, all thanks go to SPWorely for the solution. I made the following change to the above posted code and now everything is working smoothly. My fluid is no longer carbonated! :)
temp_val = atomicAdd( &dev_counters[x * sph_params.grid_y_range * sph_params.grid_z_range + y * sph_params.grid_z_range + z], 1);
dev_grid[ x * sph_params.grid_y_range * sph_params.grid_z_range * sph_params.max_cell + y * sph_params.grid_z_range * sph_params.max_cell + z * sph_params.max_cell + temp_val ] = dev_particles[tid].id;