CUDA의 블록 간 차단

Question

CUDA에서 블록 간 차단을 구현하려고하지만 심각한 문제가 발생합니다.

왜 작동하지 않는지 알 수 없습니다. 다음

__device__ void wait() { 
     while(*count != 234124) 
      ; 
    } 

프로그램으로 나는 대기를 다시 작성하는 경우에도 사실

#include <iostream> 
#include <cstdlib> 
#include <ctime> 

#define SIZE 10000000 
#define BLOCKS 100 

using namespace std; 

struct Barrier { 
    int *count; 

    __device__ void wait() { 
     atomicSub(count, 1); 
     while(*count) 
      ; 
    } 

    Barrier() { 
     int blocks = BLOCKS; 
     cudaMalloc((void**) &count, sizeof(int)); 
     cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice); 
    } 

    ~Barrier() { 
     cudaFree(count); 
    } 
}; 


__global__ void sum(int* vec, int* cache, int *sum, Barrier barrier) 
{ 
    int tid = blockIdx.x; 

    int temp = 0; 
    while(tid < SIZE) { 
     temp += vec[tid]; 
     tid += gridDim.x; 
    } 

    cache[blockIdx.x] = temp; 

    barrier.wait(); 

    if(blockIdx.x == 0) { 
     for(int i = 0 ; i < BLOCKS; ++i) 
      *sum += cache[i]; 
    } 
} 

int main() 
{ 
    int* vec_host = (int *) malloc(SIZE * sizeof(int));  
    for(int i = 0; i < SIZE; ++i) 
     vec_host[i] = 1; 

    int *vec_dev; 
    int *sum_dev; 
    int *cache; 
    int sum_gpu = 0; 

    cudaMalloc((void**) &vec_dev, SIZE * sizeof(int)); 
    cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &sum_dev, sizeof(int)); 
    cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice); 
    cudaMalloc((void**) &cache, BLOCKS * sizeof(int)); 
    cudaMemset(cache, 0, BLOCKS * sizeof(int)); 

    Barrier barrier; 
    sum<<<BLOCKS, 1>>>(vec_dev, cache, sum_dev, barrier); 

    cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost); 

    cudaFree(vec_dev); 
    cudaFree(sum_dev); 
    cudaFree(cache); 
    free(vec_host); 
    return 0; 
} 

는() 정상적으로 종료합니다. 하지만이 경우 무한 루프가 발생할 것으로 예상됩니다.

Answer 1

CUDA에서 달성하려는 목표 (블록 간 통신/동기화)는 엄격하게 가능하지 않습니다. CUDA 프로그래밍 가이드에서는 "스레드 블록을 독립적으로 실행해야합니다. 즉, 스레드 블록을 임의의 순서, 병렬 또는 직렬로 실행할 수 있어야합니다." 이러한 제한이있는 이유는 스레드 블록 스케줄러에서 유연성을 허용하고 코드가 코어 수와 함께 비례 적으로 확장되도록하기 위함입니다. 유일하게 지원되는 블록 간 동기화 방법은 다른 커널을 시작하는 것입니다. 즉 커널 시작 (동일한 스트림 내에서)은 암시 적 동기화 포인트입니다.

코드는 블록의 독립성 규칙을 위반합니다. 이는 커널의 스레드 블록이 동시에 실행된다는 것을 암시 적으로 가정하기 때문입니다 (병렬로 참조). 그러나 그들이하는 것은 아닙니다. 이것이 당신의 코드에 중요한 이유를 알기 위해 하나의 핵심만을 가진 가상의 GPU를 생각해 봅시다. 또한 두 개의 스레드 블록 만 시작한다고 가정합니다. 이 상황에서 spinloop 커널은 실제로 교착 상태가됩니다. 스레드 블록 0이 코어에서 먼저 스케줄 된 경우, 스레드 블록 1은 카운터를 갱신 할 기회가 없으므로 장벽에 도달하면 영원히 반복됩니다. 스레드 블록 0은 절대로 스왑 아웃되지 않으므로 (스레드 블록이 완료 될 때까지 실행 됨) 회전하는 동안 코어의 스레드 블록 블록 하나가 사용되지 않습니다.

일부 사람들은 당신 같은 계획을 시도해 보았습니다. 그리고 스케줄러가 예상치 못한 방법으로 블록을 일정에 세우므로 우연히 성공한 것으로 나타났습니다. 예를 들어, GPU처럼 많은 스레드 블록을 실행하면 SM이 블록이 실제로 동시에 실행된다는 것을 의미합니다. 그러나 드라이버 나 CUDA 런타임 또는 GPU를 변경하여 해당 가정을 무효화하고 코드를 위반하면 실망했습니다.

응용 프로그램의 경우 블록 간 동기화에 의존하지 않는 솔루션을 찾으십시오. CUDA 프로그래밍 모델의 의미가 바뀌지 않는 한 불가능하기 때문에 불가능합니다.

Answer 2

컴파일러 최적화 문제가 발생합니다. 나는 PTX-코드를 읽기 좋은 아니지만, 그것은 (-O0로 컴파일 된 경우에도) 컴파일러 모두에서 while -loop를 생략 한 다음과 같습니다

는 CPU 코드의 경우

.loc 3 41 0 
cvt.u64.u32  %rd7, %ctaid.x; // Save blockIdx.x to rd7 
ld.param.u64 %rd8, [__cudaparm__Z3sumPiS_S_7Barrier_cache]; 
mov.s32  %r8, %ctaid.x; // Now calculate ouput address 
mul.wide.u32 %rd9, %r8, 4; 
add.u64  %rd10, %rd8, %rd9; 
st.global.s32 [%rd10+0], %r5; // Store result to cache[blockIdx.x] 
.loc 17 128 0 
ld.param.u64 %rd11, [__cudaparm__Z3sumPiS_S_7Barrier_barrier+0]; // Get *count to rd11 
mov.s32  %r9, -1; // put -1 to r9 
atom.global.add.s32  %r10, [%rd11], %r9; // Do AtomicSub, storing the result to r10 (will be unused) 
cvt.u32.u64  %r11, %rd7; // Put blockIdx.x saved in rd7 to r11 
mov.u32  %r12, 0; // Put 0 to r12 
setp.ne.u32  %p3, %r11, %r12; // if(blockIdx.x == 0) 
@%p3 bra $Lt_0_5122; 
ld.param.u64 %rd12, [__cudaparm__Z3sumPiS_S_7Barrier_sum]; 
ld.global.s32 %r13, [%rd12+0]; 
mov.s64  %rd13, %rd8; 
mov.s32  %r14, 0; 

, 그런 행동을 방지 할 수있다 변수를 volatile 접두어로 선언하십시오. 그러나 우리는 선언 할 경우에도 countint __device__ count으로 (적절하게 코드를 변경) volatile 지정자를 추가, 그냥 내가 CUDA SDK에서 threadFenceReduction 예를보고 제안 (오류가 argument of type "volatile int *" is incompatible with parameter of type "void *"을 loke로)

컴파일을 나누기. 거기에서 그들은 거의 같은 일을하고 있지만 최종 합계를 수행 할 블록은 미리 정의 된 것이 아니라 런타임에서 선택되며 while -loop은 제거됩니다. 왜냐하면 전역 변수의 스핀 잠금이 으로 느려야하기 때문입니다.

Answer 3

블록 동기화가 가능합니다. 이 paper을 참조하십시오.
이 문서는 작동 방식에 대해 자세히 설명하지는 않지만 __syncthreads()의 작동에 의존합니다. 다른 블록이 동기화 지점에 도달하기를 기다리는 동안 현재 블록에 대한 일시 중지 장벽을 만들 수 있습니다.

이 논문에서 언급되지 않은 한 가지 항목은 블록 수가 충분히 적거나 SM의 수가 작업중인 작업에 충분히 큰 경우에만 동기화가 가능하다는 것입니다. 즉 SM이 4 개이고 5 개의 블록을 동기화하려는 경우 커널이 교착 상태가됩니다.

그들의 접근 방식을 통해 나는 여러 블록에 긴 직렬 작업을 분산시킬 수 있었고 단일 블록 접근 방식보다 30 %의 시간을 쉽게 절약 할 수있었습니다. 즉 블록 동기화는 나를 위해 일했습니다.