Вам нужно построить свою элементарную функцию, используя атомарные функции, доступные в CUDA.Я предлагаю проверить Руководство по программированию CUDA для примеров того, как построить ваши атомарные функции.Взяв приведенный там пример в качестве отправной точки, я думаю, что следующая функция делает то, что вы хотите сделать
__device__ int atomicMyFunc(int* address, int compare, int val1, int val2) {
int old = *address;
int assumed;
do {
assumed = old;
old = atomicCAS(address, assumed, assumed > compare ? val1: val2);
}
while (assumed != old);
return old;
}
Она в основном использует atomicCAS()
в цикле, чтобы сохранить результат операции только тогда, когда она имеетвычислено с использованием того же значения, которое в данный момент хранится по указанному адресу.