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

한 가지 방법으로 연대순으로 실행할을 여기

는 예제 코드입니다 이것을 달성하기 위해 Lock을 수정하려고 생각했습니다 :

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 ... 아마 누락 뭔가를 분명 포장 마차.

누구든지 도와 주시면 감사하겠습니다.

감사합니다 !!!

Answer 1

코드가 약간 변경되었습니다. 이 mutex alraedy의 값과 동일한 경우가

#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 : 이제, 원하는 출력을 생성한다. 예인 경우 0xFFFFFFFF을 mutex에 넣고 스레드가 잠금을 획득했음을 나타냅니다. mutex은 생성자에서 0으로 초기화되므로 블록 ID가 0 인 스레드 만 잠금 획득에 성공합니다. unlock에서는 원하는 블록 순서를 보장하기 위해 다음 블록 ID 색인을 mutex에 넣습니다. 또한 printf()을 CUDA 커널에 사용했기 때문에 threadfence_system()으로 전화하면 으로 전화가 걸려야합니다. 동일한 순서로 출력되는을 참조하십시오.