CUDA MUTEX How can you create mutex in Cuda?

Hi, i have a question.

How can i create mutex or Lock in CUDA without atomic functions?

[codebox]gloabal func()

{

 lock(x)

{

     x += threadIdx.x;

}

}[/codebox]

“x” in global memory

Thanks