Байт атомарной операции CUDA, чтобы заставить действовать только один поток

Question

Я пишу программу CUDA, в которой массив определен в общей памяти. Что мне нужно сделать, это разрешить только одному потоку писать каждый индекс в этом массиве, т.е. е. первый поток, достигший этой инструкции записи, должен изменить свое значение, но любые другие потоки, находящиеся в той же деформации или в следующих деформациях, должны прочитать записанное значение.

Вот фрагмент кода:

char* seq_copied = seqs + (njobNew * halfLength); //this is the shared memory array
if (seq_copied[seq_1_index] == false) { //here is the condition that I need to check with only one thread
    seq_copied[seq_1_index] = true; //and this is the write that should be written by only one thread
    printf("copy seq_shared seq_1_index = %d,  block = %d \n", seq_1_index, blockIdx.x);
}

В настоящее время происходит то, что все потоки в деформации выполняют эту точную последовательность инструкций, таким образом, оставшийся код в условии if выполняется 32 раза. Мне нужно выполнить это только один раз.

Как мне этого добиться?

Answer 1

Для этого вы можете использовать atomicCAS(). Он выполняет атомарную операцию сравнения и замены.

Эта функция будет проверять переменную и, если она соответствует определенному условию (скажем, false), она заменит ее другим значением (скажем, true). Он будет делать все эти вещи атомарно, то есть без возможности прерывания.

Возвращаемое значение атомарной функции дает нам полезную информацию в этом случае. Если возвращаемое значение равно false для приведенного выше примера, то мы можем быть уверены, что оно было заменено на true. Мы также можем быть уверены, что мы были «первым» потоком, выполнившим это условие, и все другие потоки, выполняющие аналогичную операцию, будут иметь возвращаемое значение true, а не false.

Вот рабочий пример:

$ cat t327.cu
#include <stdio.h>

__global__ void k(){

  __shared__ int flag;
  if (threadIdx.x == 0) flag = 0;
  __syncthreads();

  int retval = atomicCAS(&flag, 0, 1);
  printf("thread %d saw flag as %d\n", threadIdx.x, retval);
  // could do if statement on retval here
}


int main(){

  k<<<1,32>>>();
  cudaDeviceSynchronize();
}
$ nvcc -o t327 t327.cu
$ cuda-memcheck ./t327
========= CUDA-MEMCHECK
thread 0 saw flag as 0
thread 1 saw flag as 1
thread 2 saw flag as 1
thread 3 saw flag as 1
thread 4 saw flag as 1
thread 5 saw flag as 1
thread 6 saw flag as 1
thread 7 saw flag as 1
thread 8 saw flag as 1
thread 9 saw flag as 1
thread 10 saw flag as 1
thread 11 saw flag as 1
thread 12 saw flag as 1
thread 13 saw flag as 1
thread 14 saw flag as 1
thread 15 saw flag as 1
thread 16 saw flag as 1
thread 17 saw flag as 1
thread 18 saw flag as 1
thread 19 saw flag as 1
thread 20 saw flag as 1
thread 21 saw flag as 1
thread 22 saw flag as 1
thread 23 saw flag as 1
thread 24 saw flag as 1
thread 25 saw flag as 1
thread 26 saw flag as 1
thread 27 saw flag as 1
thread 28 saw flag as 1
thread 29 saw flag as 1
thread 30 saw flag as 1
thread 31 saw flag as 1
========= ERROR SUMMARY: 0 errors
$

Отвечая на вопрос в комментариях, мы могли бы расширить его до флага размером char, создав произвольную атомарную операцию, смоделированную после функции double atomicAdd(), указанной в руководства по программированию . Основная идея заключается в том, что мы будем выполнять atomicCAS, используя поддерживаемый размер данных (например, unsigned), и преобразуем необходимую операцию для эффективной поддержки размера char. Это делается путем преобразования адреса char в адрес с соответствующим выравниванием unsigned, а затем смещения количества char для выравнивания в соответствующей позиции байта в значении unsigned.

Вот рабочий пример:

$ cat t327.cu
#include <stdio.h>
__device__ char my_char_atomicCAS(char *addr, char cmp, char val){
  unsigned *al_addr = reinterpret_cast<unsigned *> (((unsigned long long)addr) & (0xFFFFFFFFFFFFFFFCULL));
  unsigned al_offset = ((unsigned)(((unsigned long long)addr) & 3)) * 8;
  unsigned mask = 0xFFU;
  mask <<= al_offset;
  mask = ~mask;
  unsigned sval = val;
  sval <<= al_offset;
  unsigned old = *al_addr, assumed, setval;
  do {
        assumed = old;
        setval = assumed & mask;
        setval |= sval;
        old = atomicCAS(al_addr, assumed, setval);
    } while (assumed != old);
  return (char) ((assumed >> al_offset) & 0xFFU);
}

__global__ void k(){

  __shared__ char flag[1024];
  flag[threadIdx.x] = 0;
  __syncthreads();

  int retval = my_char_atomicCAS(flag+(threadIdx.x>>1), 0, 1);
  printf("thread %d saw flag as %d\n", threadIdx.x, retval);
}


int main(){
  k<<<1,32>>>();
  cudaDeviceSynchronize();
}
$ nvcc -o t327 t327.cu
$ cuda-memcheck ./t327
========= CUDA-MEMCHECK
thread 0 saw flag as 0
thread 1 saw flag as 1
thread 2 saw flag as 0
thread 3 saw flag as 1
thread 4 saw flag as 0
thread 5 saw flag as 1
thread 6 saw flag as 0
thread 7 saw flag as 1
thread 8 saw flag as 0
thread 9 saw flag as 1
thread 10 saw flag as 0
thread 11 saw flag as 1
thread 12 saw flag as 0
thread 13 saw flag as 1
thread 14 saw flag as 0
thread 15 saw flag as 1
thread 16 saw flag as 0
thread 17 saw flag as 1
thread 18 saw flag as 0
thread 19 saw flag as 1
thread 20 saw flag as 0
thread 21 saw flag as 1
thread 22 saw flag as 0
thread 23 saw flag as 1
thread 24 saw flag as 0
thread 25 saw flag as 1
thread 26 saw flag as 0
thread 27 saw flag as 1
thread 28 saw flag as 0
thread 29 saw flag as 1
thread 30 saw flag as 0
thread 31 saw flag as 1
========= ERROR SUMMARY: 0 errors
$

Выше представлен обобщенный atomicCAS для char размера. Это позволит вам поменять любое значение char на любое другое значение char. В вашем конкретном случае, если вам нужен только логический флаг, вы можете сделать эту операцию более эффективной, используя atomicOr, как уже упоминалось в комментариях. Использование atomicOr позволит вам устранить цикл в пользовательской атомарной функции, описанной выше. Вот рабочий пример:

$ cat t327.cu
#include <stdio.h>
__device__ char my_char_atomic_flag(char *addr){
  unsigned *al_addr = reinterpret_cast<unsigned *> (((unsigned long long)addr) & (0xFFFFFFFFFFFFFFFCULL));
  unsigned al_offset = ((unsigned)(((unsigned long long)addr) & 3)) * 8;
  unsigned my_bit = 1U << al_offset;
  return (char) ((atomicOr(al_addr, my_bit) >> al_offset) & 0xFFU);
}

__global__ void k(){

  __shared__ char flag[1024];
  flag[threadIdx.x] = 0;
  __syncthreads();

  int retval = my_char_atomic_flag(flag+(threadIdx.x>>1));
  printf("thread %d saw flag as %d\n", threadIdx.x, retval);
}


int main(){
  k<<<1,32>>>();
  cudaDeviceSynchronize();
}
$ nvcc -o t327 t327.cu
$ cuda-memcheck ./t327
========= CUDA-MEMCHECK
thread 0 saw flag as 0
thread 1 saw flag as 1
thread 2 saw flag as 0
thread 3 saw flag as 1
thread 4 saw flag as 0
thread 5 saw flag as 1
thread 6 saw flag as 0
thread 7 saw flag as 1
thread 8 saw flag as 0
thread 9 saw flag as 1
thread 10 saw flag as 0
thread 11 saw flag as 1
thread 12 saw flag as 0
thread 13 saw flag as 1
thread 14 saw flag as 0
thread 15 saw flag as 1
thread 16 saw flag as 0
thread 17 saw flag as 1
thread 18 saw flag as 0
thread 19 saw flag as 1
thread 20 saw flag as 0
thread 21 saw flag as 1
thread 22 saw flag as 0
thread 23 saw flag as 1
thread 24 saw flag as 0
thread 25 saw flag as 1
thread 26 saw flag as 0
thread 27 saw flag as 1
thread 28 saw flag as 0
thread 29 saw flag as 1
thread 30 saw flag as 0
thread 31 saw flag as 1
========= ERROR SUMMARY: 0 errors
$

Эти char атомарные методы предполагают, что вы выделили массив char, размер которого кратен 4. Было бы недопустимо делать это с char массивом размера 3 (и только 3 потока) Например.