cuda
블록 간 통신

비고

마지막 블록 가드

일부 작업, 예를 들어 병행 축소로 작업하는 표를 고려하십시오. 처음에는 각 블록이 독립적으로 작업을 수행하여 일부 결과가 나옵니다. 그러나 결국에는 부분 결과를 결합하고 병합해야합니다. 전형적인 예는 큰 데이터에 대한 감소 알고리즘입니다.

일반적인 접근법은 두 개의 커널을 호출하는 것입니다. 하나는 부분 계산을위한 것이고 다른 하나는 병합을위한 것입니다. 그러나 병합을 단일 블록으로 효율적으로 수행 할 수있는 경우 커널 호출은 하나만 필요합니다. 이것은 다음과 같이 정의 된 lastBlock 가드에 의해 달성됩니다 :

__device__ bool lastBlock(int* counter) {
  __threadfence(); //ensure that partial result is visible by all blocks
  int last = 0;
  if (threadIdx.x == 0)
    last = atomicAdd(counter, 1);
  return __syncthreads_or(last == gridDim.x-1);
}

__device__ bool lastBlock(int* counter) {
  __shared__ int last;
  __threadfence(); //ensure that partial result is visible by all blocks
  if (threadIdx.x == 0) {
    last = atomicAdd(counter, 1);
  }
  __syncthreads();
  return last == gridDim.x-1;
}

이러한 가드를 사용하면 마지막 블록은 다른 모든 블록에서 생성 된 모든 결과를 볼 수 있으며 병합을 수행 할 수 있습니다.

__device__ void computePartial(T* out) { ... }
__device__ void merge(T* partialResults, T* out) { ... }

__global__ void kernel(int* counter, T* partialResults, T* finalResult) {
    computePartial(&partialResults[blockIdx.x]);
    if (lastBlock(counter)) {
      //this is executed by all threads of the last block only
      merge(partialResults,finalResult);
    }
}

가정 :

• 카운터는 전역 메모리 포인터 여야하며 커널이 호출 되기 전에 0으로 초기화됩니다.
• lastBlock 함수는 모든 블록의 모든 스레드에 의해 균등하게 호출됩니다.
• 커널은 1 차원 격자로 호출됩니다 (예제의 단순함을 위해).
• T 원하는 모든 유형의 이름을 지정하지만 예제는 C ++의 템플리트가 아니며

글로벌 작업 대기열

일련의 작업 항목을 고려하십시오. 각 작업 항목을 완료하는 데 필요한 시간은 크게 다릅니다. 블록 간의 작업 분배를 균형 잡기 위해 각 블록이 이전 항목이 완료 될 때만 다음 항목을 가져 오는 것이 현명합니다. 이는 선험적으로 블록에 항목을 할당하는 것과는 대조적입니다.

class WorkQueue {
private:
  WorkItem* gItems;
  size_t totalSize;
  size_t current;
public:
  __device__ WorkItem& fetch() {
    __shared__ WorkItem item;
    if (threadIdx.x == 0) {
      size_t itemIdx = atomicAdd(current,1);
      if (itemIdx<totalSize)
        item = gItems[itemIdx];
      else
        item = WorkItem::none();
    }
    __syncthreads();
    return item; //returning reference to smem - ok
  }
}

가정 :

• gitem 배열뿐만 아니라 WorkQueue 객체는 전역 메모리에 상주합니다.
• 페치하는 커널의 WorkQueue 오브젝트에 새로운 작업 항목이 추가되지 않습니다.
• WorkItem 은 작업 할당의 작은 표현입니다 (예 : 다른 객체에 대한 포인터).
• WorkItem::none() 정적 멤버 함수는 작업을 전혀 나타내지 않는 WorkItem 객체를 만듭니다.
• WorkQueue::fetch() 는 블록의 모든 스레드에 의해 균일하게 호출되어야합니다.
• 그 사이에 또 ​​다른 __syncthreads() 가 없으면 WorkQueue::fetch() 를 호출하지 않습니다. 그렇지 않으면 경쟁 조건이 나타납니다!

이 예제에는 WorkQueue 초기화하거나 채우는 방법은 포함되어 있지 않습니다. 그것은 다른 커널 또는 CPU 코드에 의해 수행되며 꽤 간단합니다.