쉐이더 유닛이 지수를 계산합니까?

Question

http://us.hardware.info/reviews/5419/nvidia-geforce-gtx-titan-z-sli-review-incl-tones-tizair-system 에 "GTX Titan-Z"에는 5760 개의 쉐이더가 있다고 표시됩니다. 또한 여기에 "GTX Titan-Z"에는 2x GK110 GPU가 있습니다.

CUDA exp() expf() and __expf()은 cuda에서 지수를 계산할 수 있음을 언급합니다.

내가 500,000 000 (500 억)의 복식 배열을 가지고 있다고 가정 해 봅시다. 배열의 각 값의 지수를 계산하고 싶습니다. 무엇을 기대해야하는지 : 5760 개의 쉐이더 유닛이 exp를 계산할 수 있습니까? 아니면이 작업은 2 개의 GK110 GPU에서만 수행 할 수 있습니까? 성능 차이가 심하기 때문에 CUDA로 응용 프로그램을 다시 작성하면 속도가 느려질 수 있습니다.

즉, 57000 개의 스레드를 만들어 500 000 000 지수를 계산할 수 있습니까?

Answer 1

GTX Titan Z는 듀얼 GPU 장치입니다. 카드의 두 GK110 GPU는 각각 384 비트 메모리 인터페이스를 통해 자체 6GB의 고속 메모리에 연결됩니다. 각 메모리의 이론적 대역폭은 336GB/초입니다. GTX Titan Z에서 사용 된 특정 GK110 변형은 SMX라고 불리는 15 개의 실행 단위 클러스터로 구성됩니다. 각 SMX는 차례대로 192 개의 단 정밀도 부동 소수점 단위, 64 개의 배정 밀도 부동 소수점 단위 및 기타 여러 단위로 구성됩니다.

GK110의 각 배 정밀도 단위는 클럭주기 당 하나의 FMA (퓨즈 된 다중 덧셈) 또는 FMUL 하나 또는 FADD를 실행할 수 있습니다. 705 MHz의 기본 클럭에서 Titan Z의 초당 GK110 GPU가 실행할 수있는 DP 작업의 최대 총 수는 705e6 * 15 * 64 = 676.8e9입니다. 모든 연산이 FMA라고 가정 할 때, 이는 1.3536 배정 밀도 TFLOPS와 같습니다. 카드가 2 개의 GPU를 사용하기 때문에 GTX Titan Z의 총 DP 성능은 2.7072 TFLOPS입니다.

GPU는 CPU와 마찬가지로 다양한 정수 및 부동 소수점 단위로 범용 계산을 제공합니다. 또한 GPU는 역수, 역수 제곱근, 사인 곡선, 코사인, 지수 기수 2와 같은 자주 사용되는 함수에 대한 대략적인 단 정밀도 근사를 계산할 수있는 특수 함수 유닛 (MUFU = mu lti fu이라는 단위를 GK110에 제공함) , 그리고 대수를 기반으로합니다. 지수 연산과 관련하여 표준 단 정밀도 연산 함수 exp2f()은 MUFU 명령어 (MUFU.EX2)에 어느 정도 직접 매핑하는 유일한 함수입니다. 하드웨어가 특수 기능 장치에서 비정규 피연산자를 지원하지 않기 때문에 컴파일 모드에 따라이 하드웨어 명령어 주위에 얇은 래퍼가 있습니다.

CUDA의 다른 모든 지수는 소프트웨어 서브 루틴을 통해 수행됩니다. 표준 단일 정밀도 함수 expf()은 하드웨어의 exp2 기능을 상당히 무거운 래퍼입니다. 배정도 exp() 함수는 미니 맥스 다항식 근사법을 기반으로하는 순수한 소프트웨어 루틴입니다. 그것의 완전한 소스 코드는 CUDA 헤더 파일 math_functions_dbl_ptx3.h (CUDA 6.5에서 해당 파일의 1706 번째 라인에서 DP exp() 코드 시작)에 표시됩니다. 보시다시피, 계산에는 주로 정수 및 일부 단 정밀도 부동 소수점 연산뿐만 아니라 배정 밀도 부동 소수점 연산이 포함됩니다. 을 호출하는 이진 실행 파일을 역 어셈블하여 cuobjdump --dump-sass으로 기계 코드를 볼 수도 있습니다.

성능면에서 CUDA 6.5에서 배정도 exp() 함수는 Tesla K20 (1.170 DP TFLOPS)에서 초당 25e9 함수 호출의 처리량을 나타냅니다. DP exp()을 호출 할 때마다 8 바이트 소스 피연산자를 사용하고 8 바이트 결과를 생성하므로 약 400GB/초의 메모리 대역폭과 같습니다.Titan Z의 각 GK110은 Tesla K20의 GK110보다 15 % 더 높은 성능을 제공하기 때문에 처리량과 대역폭 요구 사항이 그에 따라 증가합니다. 필요한 대역폭이 GPU의 이론적 인 메모리 대역폭을 초과하기 때문에 단순히 배열에 DP exp()을 적용하는 코드는 메모리 대역폭에 완전히 바인딩됩니다.

GPU의 기능 단위 수와 실행중인 스레드 수는 처리 할 수있는 배열 요소 수와는 관계가 없지만 이러한 처리 성능에 영향을 미칠 수 있습니다. 스레드에 대한 배열 요소의 매핑은 프로그래머가 자유롭게 선택할 수 있습니다. 한 번에 처리 할 수있는 배열 요소의 수는 GPU 메모리 크기의 함수입니다. CUDA 소프트웨어 스택은 일반적으로 100MB 정도의 메모리를 필요로하기 때문에 장치의 모든 원시 메모리를 사용자 코드에 사용할 수있는 것은 아닙니다. 어레이에 DP exp()을인가하기위한 예시적인 맵핑이 코드에 나타내

는

__global__ void exp_kernel (const double * __restrict__ src, 
          double * __restrict__ dst, int len) 
{ 
    int stride = gridDim.x * blockDim.x; 
    int tid = blockDim.x * blockIdx.x + threadIdx.x; 
    for (int i = tid; i < len; i += stride) { 
     dst[i] = exp (src[i]); 
    } 
}  

#define ARRAY_LENGTH (500000000) 
#define THREADS_PER_BLOCK (256) 
int main (void) { 
    // ... 
    int len = ARRAY_LENGTH; 
    dim3 dimBlock(THREADS_PER_BLOCK); 
    int threadBlocks = (len + (dimBlock.x - 1))/dimBlock.x; 
    if (threadBlocks > 65520) threadBlocks = 65520; 
    dim3 dimGrid(threadBlocks); 
    double *d_a = 0, *d_b = 0; 

    cudaMalloc((void**)&d_a, sizeof(d_a[0]), len); 
    cudaMalloc((void**)&d_b, sizeof(d_b[0]), len); 
    // ... 
    exp_kernel<<<dimGrid,dimBlock>>>(d_a, d_b, len); 
    // ... 
}