How to implement spin lock / problem with atomic operations

Accelerated Computing CUDA CUDA Programming and Performance
1

Hey there,

currently I’m trying to implement a spin lock to manage access on local memory. As far as I understand the whole thing the following code should just do that:

#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable

function()

{

local_lock_p[0] is initialized to 0 and is located in local memory

//try to aquire a spin lock

  while(atom_xchg(local_lock_p,1,1) ) {};

[...do something]

//return lock

  atom_xchg(local_lock_p,1,0);

}

But now I’m getting the following error:

sm_11 does not support intrinsic %llvm.ptx.atom.cmpxchg.l.ii

How else am I supposed to implement the spin lock (or any other lock) if I can’t use atomic operations?

Driver is 197.59, GPU is Quadro FX 3700. Is this solved with a newer driver version, I’d prefer to keep the old one atm since I’m sharing the workstation with others.

Cheers

Magnus

2

I’m still looking for a solution to this problem!

3

Ok, since nobody replied I had no chance but try out the newest driver (266.45) and now I get the following error:

My function head looks like that:

__kernel void findLocalMax(

				  const __global float * lower_layer_p 

				, const __global float * middle_layer_p

				, const __global float * upper_layer_p

				, __global	 float4 * result  //4 times layer size

				, const	  uint   width_layer

				, const	  uint   height_layer

				, __local int  * local_lock_p

				, __local float * local_mem_p

				)
4

I tried to get more information regarding alignement. I found the following functions for querrying the device about alignement information.

cl_uint info = 0;

clGetDeviceInfo(l_devices,CL_DEVICE_MEM_BASE_ADDR_ALIGN,sizeof(cl_uint), (void*)&info, NULL);

	

std::cout << "CL_DEVICE_MEM_BASE_ADDR_ALIGN " << info << std::endl;

info = 0;

clGetDeviceInfo(l_devices,CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(cl_uint), (void*)&info, NULL); 

std::cout << "CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE " << info << std::endl;

The result was:

That sounds quite high, or do I misunderstand the function?

5

Hello,

I have the same Problem. First i have got the same error

and after a driver update (to 270.81) i get now again the same error like maggnet

Has nobody an solution or a hint?

I think my GPU (GeForce 9800 GTX/9800 GTX+) dont support the local atomic extensions, but iam not sure.

6

Try using clGetDeviceInfo with CL_DEVICE_EXTENSIONS to find if your GPU supports this.