Cuda Atomic Lock: потоки в последовательности

Question

У меня есть код, раздел которого должен быть выполнен критически. Я использую блокировку для этого фрагмента кода, чтобы каждый поток ядра (настроенный с одним потоком на блок) выполнял этот фрагмент кода атомарно. Меня беспокоит порядок потоков - мне нужно, чтобы потоки выполнялись в хронологическом порядке в соответствии с их индексами (или фактически в порядке их blockIdx), от 0 до 10 (вместо случайного, например, 5, 8, 3, 0, ... и т. Д.) Возможно ли это сделать?

Вот пример кода:

#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<math_functions.h>
#include<time.h>
#include<cuda.h>
#include<cuda_runtime.h>

// number of blocks
#define nob 10

struct Lock{
  int *mutex;
  Lock(void){
    int state = 0;
    cudaMalloc((void**) &mutex, sizeof(int));
    cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
  }
  ~Lock(void){
    cudaFree(mutex);
  }
  __device__ void lock(void){
    while(atomicCAS(mutex, 0, 1) != 0);
  }
  __device__ void unlock(void){
    atomicExch(mutex, 0);
  }
};


__global__ void theKernel(Lock myLock){
  int index = blockIdx.x; //using only one thread per block

  // execute some parallel code

  // critical section of code (thread with index=0 needs to start, followed by index=1, etc.)
  myLock.lock();

  printf("Thread with index=%i inside critical section now...\n", index);

  myLock.unlock();
}

int main(void)
{
  Lock myLock;
  theKernel<<<nob, 1>>>(myLock);
  return 0;
}

что дает следующие результаты:

Thread with index=1 inside critical section now...
Thread with index=0 inside critical section now...                                                                                                                                   
Thread with index=5 inside critical section now...                                                                                                                                            
Thread with index=9 inside critical section now...
Thread with index=7 inside critical section now...
Thread with index=6 inside critical section now...
Thread with index=3 inside critical section now...
Thread with index=2 inside critical section now...
Thread with index=8 inside critical section now...
Thread with index=4 inside critical section now...

Я хочу, чтобы эти индексы начинались с 0 и выполнялись в хронологическом порядке до 9.

Один из способов, с помощью которого я решил изменить блокировку, заключается в следующем:

struct Lock{
  int *indexAllow;
  Lock(void){
    int startVal = 0;
    cudaMalloc((void**) &indexAllow, sizeof(int));
    cudaMemcpy(indexAllow, &startVal, sizeof(int), cudaMemcpyHostToDevice);
  }
  ~Lock(void){
    cudaFree(indexAllow);
  }
  __device__ void lock(int index){
    while(index!=*indexAllow);
  }
  __device__ void unlock(void){
    atomicAdd(indexAllow,1);
  }
};

а затем просто инициализировать блокировку, передав индекс в качестве аргумента:

myLock.lock(index);

но это останавливает мой компьютер... Я, вероятно, упускаю что-то очевидное.

Если кто-то может помочь, я буду благодарен!

Спасибо!!!

Answer 1

Я немного изменил твой код. Теперь он производит желаемый результат:

#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<math_functions.h>
#include<time.h>
#include<cuda.h>
#include<cuda_runtime.h>

// number of blocks
#define nob 10

struct Lock{
  int *mutex;
  Lock(void){
    int state = 0;
    cudaMalloc((void**) &mutex, sizeof(int));
    cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
  }
  ~Lock(void){
    cudaFree(mutex);
  }
  __device__ void lock(uint compare){
    while(atomicCAS(mutex, compare, 0xFFFFFFFF) != compare);    //0xFFFFFFFF is just a very large number. The point is no block index can be this big (currently).
  }
  __device__ void unlock(uint val){
    atomicExch(mutex, val+1);
  }
};


__global__ void theKernel(Lock myLock){
  int index = blockIdx.x; //using only one thread per block

  // execute some parallel code

  // critical section of code (thread with index=0 needs to start, followed by index=1, etc.)
  myLock.lock(index);
  printf("Thread with index=%i inside critical section now...\n", index);
  __threadfence_system();   // For the printf. I'm not sure __threadfence_system() can guarantee the order for calls to printf().
  myLock.unlock(index);
}

int main(void)
{
  Lock myLock;
  theKernel<<<nob, 1>>>(myLock);
  return 0;
}

lock() функция принимает compare в качестве параметра и проверяет, равно ли оно значению alraedy в mutex, Если да, это ставит 0xFFFFFFFF в mutex чтобы указать, что блокировка получена потоком. Поскольку mutex инициализируется в конструкторе 0, только поток с идентификатором блока 0 будет успешен в получении блокировки. в unlockмы помещаем индекс следующего блока в mutex чтобы гарантировать желаемый заказ. Кроме того, потому что вы использовали printf() внутри ядра CUDA, я думаю, что вызов threadfence_system() требуется, чтобы вы видели их в выводе в том же порядке.