nvcc - 컴파일 타임의 아치에 따라 다른 블록 크기

Question

나는 최고의 성능을 보여주는 커널을 가지고 있습니다 다른 Kepler 및 Fermi 하드웨어에서 실행될 때 블록 크기.컴파일 타임에 현재 아키텍처 타겟을 확인하고 THREADS_PER_BLOCK i) 매크로를 사용하여 커널을 시작합니다.ii) 필요한 블록 수를 결정합니다.iii) 커널에서 공유 메모리 크기를 정적으로 설정합니다.

아래는 내가하려는 일을 보여줍니다.GK104 하드웨어를 목표로 삼고 있으므로 다음을 사용한다고 가정합니다. nvcc -arch=sm_30.이로 인해 여전히 THREADS_PER_BLOCK = 256 ~부터 __CUDA_ARCH__ 호스트 코드 컴파일에 대해 정의되지 않았습니다.(예를 들어 이해합니다. 이 답변, 왜 이런 식으로 작동할 수 없는가.)

#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif

__global__ void some_kernel(int* a, int* b) {
    __shared__ sm_data[THREADS_PER_BLOCK];
    // Do something.
}

int main(void) {
    // Initialize data.
    // Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
    some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
    return 0;
}

런타임에 장치 속성을 확인하고 동적 공유 메모리를 사용할 수 있지만 컴파일 타임에 예를 들어 하드 코딩할 수 있는지 알고 싶습니다.-DFERMI 또는 -DKEPLER를 수동으로 추가하고 설정해야 함 THREADS_PER_BLOCK 그것을 바탕으로.참고:이 코드를 사용하는 모든 사용자는 스스로 코드를 컴파일할 것입니다. 거의 확실히 하나의 아키텍처에 대해서는 이것이 불합리한 선택이 아닙니다.통과에 비추어 볼 때 불필요해 보입니다. -arch= 깃발.

Answer 1

nvcc 컴파일러는 로컬에서 사용 가능한 GPU를 감지하지 않으며 기본적으로 항상 SM 1.0을 대상으로 합니다.그렇지 않으면 다른 시스템에서 구축할 때 상당히 혼란스러운 동작이 발생할 수 있습니다.

사용 가능한 장치에 맞게 컴파일하려면 사용자에게 SM 버전을 지정하도록 요청하거나 빌드 시간 동안 일부 감지 코드를 실행해야 합니다.나는 하드웨어 감지 코드를 런타임에 넣은 다음 원하는 대로 커널 실행을 구성하는 것이 더 쉽다고 확신합니다.