1

セクションを批判的に実行する必要があるコードがあります。そのコードにロックを使用して、カーネルの各スレッド (ブロックごとに 1 つのスレッドで設定) がそのコードをアトミ​​ックに実行するようにしています。スレッドの順序は私を悩ませているものです-0から10までのインデックスに従って（または実際にはblockIdxの順序で）時系列でスレッドを実行する必要があります（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 まで実行するようにします。

これを達成するためにロックを変更する方法の 1 つは、次のとおりです。

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);

しかし、これは私のPCを失速させます...おそらく明らかな何かが欠けています。

誰かが助けることができれば、私はそれを感謝します!

ありがとう！！！

4

1 に答える 1

2

コードを少し変更しました。これで、目的の出力が生成されます。

#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それがすでに の値と等しいかどうかをチェックしますmutex。はいの場合、ロックがスレッドによって取得されたことを示すために に0xFFFFFFFF挿入されます。mutexはコンストラクターで 0 で初期化されるため、mutexブロック ID が 0 のスレッドのみがロックの取得に成功します。では、希望する順序を保証するためにunlock、次のブロック ID インデックスを に配置します。mutexまた、CUDA カーネル内で使用しているため、出力でそれらを同じ順序で表示するには、への呼び出しが必要printf()だと思います。threadfence_system()

于 2014-09-12T00:05:49.407 に答える