¿Cómo puedo implementar una función atómica personalizada que involucre varias variables?

Me gustaría implementar esta función atómica en CUDA:

__device__ float lowest;   // global var
__device__ int  lowIdx;    // global var
float realNum;   // thread reg var
int index;       // thread reg var

if(realNum < lowest) {
 lowest= realNum;  // the new lowest
 lowIdx= index;    // update the 'low' index
}

No creo que pueda hacer esto con ninguna de las funciones atómicas. Necesito bloquear un par de ubicaciones de memoria global para un par de instrucciones. ¿Podría implementarlo con el código PTXAS (ensamblado)?

Respuestas a la pregunta(2)

Su respuesta a la pregunta