2012-01-24 5 views
4

저는 많은 BLAS gemv 연산을 수행하는 함수를 작성하고 있습니다.OpenCL 또는 CUDA 호출의 오버 헤드가 있습니까?

GPU에서이 작업을 수행하고 싶습니다. cuBlas를 사용해 보았습니다.

제 문제는 제 행렬과 벡터가 100x100 행렬과 100 벡터로 작다는 것입니다. CuBlas는 CPU에 비해 ​​오랜 시간이 걸리며 그 이유는 CPU의 빠른 캐시와 GPU 호출에 큰 오버 헤드가 혼합 된 이유입니다.

따라서 GPU로 통화하는 데 걸리는 시간을 측정하는 현명한 방법을 알아 내려고하고 있습니다.

CUDA가 호출을 설정하고 그래픽 프로세서로 전송하는 데 걸리는 시간입니다. 실제로 매트릭스 - 벡터 곱셈을 수행하는 데 걸리는 시간은 계산되지 않습니다.

어떻게해야합니까?

+0

CUDA에 작업을 보낼지 아니면 관심을 끌지 만 동적으로 작업을 선택할 수 있도록하려면? – Rup

+0

@Rup : 전화가 실제로 들리는 것이 무엇인지 파악하는 데 관심이 있습니다. 느린 코드가 내 잘못인지 아니면 단순히 건축물의 제품인지 판단하십시오. –

+1

소량의 데이터의 경우, 이는 여러분에게 피해를주는 오버 헤드뿐만 아니라 가능한 병렬 처리의 부족도 있습니다. GPU는 대기 시간을 숨기기에 충분한 스레드를 가지고 있습니다 (CPU가 GPU에서 많이 나빠집니다). 호출 오버 헤드가 없어도 작업이 ** 많은 스레드 **로 분할되어 있지 않으면 GPU가 CPU보다 느려질 수 있습니다. 많은 것은 쉽게 수천 개의 스레드를 의미 할 수 있습니다. – Grizzly

답변

8

업데이트 다음 결과 (엔비디아 7800 GTX) 2005 하드웨어에 손으로 쓴 FFT의 GPU의 알고리즘하지만, CPU-GPU의 전송하세요의 원리를 보여준다는 오버 헤드가없는

병목 GPU 프로그램을 컴파일하고 GPU와 호스트간에 데이터를 전송합니다. CPU는 캐시에서 완전히 수행 할 수있는 기능에 고도로 최적화되어 있으며 DDR3 메모리의 대기 시간은 GPU를 서비스하는 PCI-Express 버스보다 훨씬 낮습니다. GPU FFT 루틴 (CUDA 이전)을 작성할 때이 점을 직접 경험했습니다. this related question을 참조하십시오.

 
N  FFTw (s) GPUFFT (s) GPUFFT MFLOPS GPUFFT Speedup 
8  0   0.00006  3.352705  0.006881 
16  0.000001 0.000065 7.882117  0.010217 
32  0.000001 0.000075 17.10887  0.014695 
64  0.000002 0.000085 36.080118  0.026744 
128  0.000004 0.000093 76.724324  0.040122 
256  0.000007 0.000107 153.739856  0.066754 
512  0.000015 0.000115 320.200892  0.134614 
1024 0.000034 0.000125 657.735381  0.270512 
2048 0.000076 0.000156 1155.151507  0.484331 
4096 0.000173 0.000215 1834.212989  0.804558 
8192 0.000483 0.00032  2664.042421  1.510011 
16384 0.001363 0.000605 3035.4551  2.255411 
32768 0.003168 0.00114  3450.455808  2.780041 
65536 0.008694 0.002464 3404.628083  3.528726 
131072 0.015363 0.005027 3545.850483  3.05604 
262144 0.033223 0.012513 3016.885246  2.655183 
524288 0.072918 0.025879 3079.443664  2.817667 
1048576 0.173043 0.076537 2192.056517  2.260904 
2097152 0.331553 0.157427 2238.01491  2.106081 
4194304 0.801544 0.430518 1715.573229  1.861814 

위의 표는 커널 크기에 따라 GPU FFT 구현과 CPU 구현의 타이밍을 보여줍니다. 크기가 더 작 으면 GPU에서/GPU 로의 데이터 전송이 우위를 차지합니다. 더 작은 커널은 CPU에서 수행 될 수 있으며 일부 구현/크기는 전체 캐시에서 수행 될 수 있습니다. 따라서 CPU를 소규모 작업에 최적의 선택으로 만듭니다.

반면에 GPU에서/로의 최소한의 이동으로 큰 작업 일괄 작업을 수행해야한다면 GPU가 CPU 손을 이길 것입니다.

예제에서 효과를 측정하는 한, 위와 같은 실험을 수행하는 것이 좋습니다. 행렬의 각 크기에 대해 계산 된 FLOPS를 계산하고 다양한 크기의 행렬에 대해 CPU 및 GPU에서 테스트를 실행하십시오. CSV 파일에 GPU 대 CPU의 크기, 시간 및 FLOPS를 출력합니다. 어떤 프로파일 링이라도 코드를 수백 번 반복하고 모든 것을 수행 한 다음 반복을 통해 총 시간을 나누어 루프 시간을 확보하십시오. 알고리즘에 허용되는 경우 다른 모양의 행렬을 사용해보십시오 (예 : 100x10이 아닌 10x100).

이 데이터를 사용하면 오버 헤드가 무엇인지 느낄 수 있습니다. 똑같은 실험을 반복해서 찾으려면 GPU에서 실행 된 내부 쉐이더 코드를 아무 작업없이 대체하십시오 (단순히 입력에서 출력으로 복사). 이 도움이

희망,

1

당신은 이벤트가 대기열에 제출, 시작하고 버퍼 전송 이벤트에 clGetEventProfilingInfo를 사용하여 완성 된 장치에서 나노초의 시간을 얻을 수 있습니다.

대한 추가 정보를 원하시면, 어떻게 여기까지 설정 : http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetEventProfilingInfo.html

나는 100 × 100 행렬을 위해, 당신은 사각 사각의 CPU에 고집 더 나을 수 있다고 생각합니다. 동시에 많은 것을 배가하지 않는 한, (작은) 전송 오버 헤드와 대개 훨씬 낮은 클럭 속도로 인해 gpu의 이점은 거의 눈에 띄지 않습니다. 가능한 한 많은 로컬 데이터를 사용하기 위해 커널을 조정해야합니다. 하드웨어에 작업 그룹당 32KB가 있으며 두 개의 100x100 행렬을 보유해야합니다. 내장 된 내적 함수도 매우 편리해야합니다. 그들은 커널을 최적화에 대해 자세히 이야기 http://developer.amd.com/afds/pages/OLD/sessions.aspx 및 최적의 크기를 하드 코딩 :

는 ADFS 작년에 이에 대한 멋진 이야기가 있었다 (2908 sessionId가 참조).

1

매트릭스가 이미 GPU에 있습니까? 그렇지 않은 경우 CUBLAS는 추가 오버 헤드 인 thunking으로 알려진 메시지를 전송할 수 있습니다.

또한 작은 계산에서는 GPU가 실제로 빛나지 않습니다. 결과를 다시 전송해야하기 때문에 CPU보다 느릴 수 있습니다. 가능한 경우 더 큰 행렬을 사용하십시오. 그렇지 않으면 스트림 (cudaStream_t)을 사용하여 GPU에서 여러 개의 병렬 계산을 시작할 수 있습니다. 당신이 CUDA에서 커널의 실행 시간을 측정하려면

, 당신은 묶어야합니다 그 CUDA 런타임 API 사용하는 경우 다음과 같이 행사에 (또는 GPU에 계산 아무것도) :

cudaEvent_t start, stop; 

cudaEventRecord(&start); 

struct timeval cpuStart, cpuEnd; 

gettimeofday(&cpuStart, 0); // get start time on CPU 

// Do something with CUDA on the GPU, e.g. call kernels, transfer memory, ... 

gettimeofday(&cpuEnd, 0); // get end time on CPU 

double seconds = cpuEnd.tv_sec - cpuStart.tv_sec; 
double microseconds = cpuEnd.tv_usec - cpuStart.tv_usec; 
double cpuDuration = (seconds * 1.0e6 + microseconds)/1.0e3; // in milliseconds 

cudaEventRecord(&stop); 

// Wait until the stop event occurred 
cudaError_t eventResult; 

do 
{ 
    eventResult = cudaEventQuery(stop); 
} 
while (eventResult == cudaErrorNotReady); 

// Assert there was no error; check the CUDA Toolkit Reference for further info 
assert(cudaSuccess == eventResult); // requires #include <assert.h> or <cassert> 

// Retrieve the time 
float gpuDuration = 0.0; // in milliseconds 
cudaEventElapsedTime(&gpuDuration, start, stop); 

// Release the event objects 
cudaEventDestroy(stop); 
cudaEventDestroy(start); 

당신은 이전 호출에서 오류가 발생할 수 있습니다으로 디버깅 시간의 결과로, (적어도 어설 포함) CUDA 모든 호출의 에러 코드를 확인 할 수 있습니다 ...

(참고 : 나는 주로 사용 CUDA 드라이버 API이므로 상자에서 작동하지 않을 수 있습니다. 죄송합니다.)

EDIT : 커널의 지속 시간이 아니라 호출 자체를 측정하고자하는 것을 보았습니다. 호출에 대한 CPU 시간 만 측정하면됩니다. 위의 업데이트 된 코드를 참조하십시오. Windows (AFAIK)에서는 gettimeofday를 사용할 수 없으므로이 기능은 Linux에서만 작동합니다.

+1

Windows에서는 [QueryPerformanceCounter] (http : // msdn. microsoft.com/en-us/library/ms644904) 또는 [GetSystemTime] (http://msdn.microsoft.com/en-us/library/windows/desktop/ms725473.aspx) 등 – Rup

+0

나는 모든 기기에서 데이터를 가져오고 간단히 Ax-> y를 수행 한 다음 기기에서 y를 유지하면됩니다. –

+1

이 경우 cUBLAS는 cublasDgemm() 호출을 통해 gettimeofday() (또는 Windows에서 유사한 메소드)를 실행하여 실제 커널을 실행해야하는 시간을 측정 할 수 있습니다. 직접 해보지는 않았지만 Parallel Nsight (Windows) 또는 Visual Compute Profiler (Linux의 툴킷에 포함)를 사용할 수 있습니다. 지금은 찾을 수 없지만 CUDA 4의 프로파일 링 후크에 대해 알게되었습니다. EDIT : CUDA 프로파일 링에 대한 흥미로운 정보가있는 PDF가 있습니다. http : // bit .ly/zn6jbP –

1

호출 오버 헤드를 찾으려면 가능한 한 적게 수행하는 CUDA 커널을 호출하십시오.

for (int i=0; i<NLoops; i++) { 
    gettimeofday(&cpuStart, 0); // get start time on CPU 

    // Call minimal CUDA kernel 

    gettimeofday(&cpuEnd, 0); // get end time on CPU 

    // save elapsed time 
} 

위의 Alex P. 코드를 따르십시오.

커널에서 처리하는 횟수가 적을수록 호출 오버 헤드 만 더 많은 시간 차이가 발생합니다.

NLoops (아마도 1,000,000 개)에 대한 적절한 값을 찾으려고 조금 실험 해보십시오. 경과 시간이 타이머 간격보다 길어야합니다. 그렇지 않으면 모두 0으로 끝납니다. 이런 일이 발생하면 예측할 수있는 고정 된 시간 간격으로 실행되는 커널 코드를 작성하십시오 (각각 x 사이클의 n 루프).

cpuStart와 cpuEnd (인터럽트 처리와 같은) 사이에 발생할 수있는 모든 비 CUDA 계산을 제거하기는 어렵지만 여러 번의 실행과 평균화를 수행하면 좋은 결과를 얻을 수 있습니다.