atom_add for float

Discussion created by chevydevil on Jun 20, 2010
Latest reply on Jun 24, 2010 by chevydevil

Hello. I need to sum up some forces for a deformable body mass-spring system with volume preservation. Every workitem should add part of a force for its involved masspoints(vertices). I thought i can do this with an atomic add but in the 1.0 spec only integer and long types are allowed. So is there a way to do this or do i need another kernel for synchronisation?

Edit: Some Masspoints(vertices) are be shared by different workitems. Thats the whole thing.

Edit2: Here some source code example:

This doesn't work either. It seems i really can use only __global int or _global long for atomic operations.

#pragma OPENCL EXTENSION cl_khr_fp64 : enable #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable typedef struct { float4 mpMass; float4 rl1, rl2; float m_v0, a, b, c; int4 mps; } tetra; __kernel void force_calc(__global int4* _force, __global tetra* _tetra) { int id = get_global_id(0); __global tetra *t = &_tetra[id]; int mpid = t->mps.x; __global int4 *f = &_force[mpid]; (void) atom_add(f->x,10); }