CUDA의 점유율 관리

Question

안녕하십니까.

나는 쿠다 프로그래밍을 배우기 시작하고 있으며 성능을 극복하려고합니다. 본인은 우리가 고려 네 가지를해야 좋은 성능을 가지고 그 CUDA 웹 사이트에서 읽어

http://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/achievedoccupancy.htm

-warps 당 SM (시스템 멀티) SM - 공유 당 SM -register 당 -blocks SM 당 메모리

그래서 첫 번째 것들로 넘어가겠습니다. 그리고 GPU에 따라 SM 당 최대 워프와 SM 당 블록 수에 따라 커널의 크기를 정의했습니다. 내 작업은 어떤 방법이 나아 졌는지 측정하기 위해 10 억 개의 합계를 수행합니다.

내가하는 일은 각 반복마다 점유를 극대화하는 for 루프를 실행하는 for 루프입니다. 엔비디아 1080 GPU에 대한 예를 들어 본인은 :이 SM 당 총 2,048 실에 제공하고 최대 인원을 보장

int max_blocks = 32; //maximum number of active blocks per SM int max_threads_per_Block = 64; //maximum number of active threads per SM int max_threads = 2048;

. 이 GPU는 각각 32 개의 스레드로 64 개의 활성 워프를 가질 수 있습니다. 이 GPU에서 하나의 활성 블록에는 2 개의 워프가 있으며 이는 각 블록이 동시에 64 개의 활성 스레드를 가질 수 있음을 의미합니다. 다음과 같이 함께 내가 커널을 실행 :

dim3 threadsPerBlock(max_threads_per_Block); dim3 numBlocks(max_blocks); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads);

내가 놀라 울 정도로 알 것은이 내가 좋아하는 바로이 커널을 실행하는 경우 :

int N = total_ops; //in this case one thousand millions dim3 threadsPerBlock(256); dim3 numBlocks(2*N/threadsPerBlock.x); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,);

성능이 좋습니다 (t ime consumed). 이상한 일을 피하기 위해 같은 실험을 5 번 반복합니다. 제 질문은 : 컴파일러와 런타임 API가하는 것보다 더 나은 결과를 얻기 위해 점유를 관리 할 수있는 방법이 있습니까? 나는 내가하려는 최적화가 이미 GPU에 의해 어떤 방식 으로든 관리되고 있음을 이해합니다. 나는 좋은 성과를 내기 위해 우리가 소프트웨어를 어떻게 시작해야 하는지를 설명하는 문서 (위의 링크)가 있다면이를 통제하는 방법이어야한다고 생각합니다. 완전히 한 SM를로드하는 데 필요한

감사

첫 번째 예에서

Answer 1

,

int max_blocks = 32;   //maximum number of active blocks per SM 
int max_threads_per_Block = 64; //maximum number of active threads per SM 
int max_threads = 2048; 

dim3 threadsPerBlock(max_threads_per_Block); 
dim3 numBlocks(max_blocks); 
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads); 

당신은 블록 당 많은 블록과 스레드를 실행하고 있습니다. 그러나 GTX 1080은 SM이므로 인원은 1/20 = 5 %입니다. 두 번째 예에서

, 100 %의 점유율을 얻기 위해 필요에 따라 GPU가 병렬로 많은 사람을 실행할 수 있도록 블록의 다수를 출시하고

int N = total_ops;    //in this case one thousand millions 
dim3 threadsPerBlock(256); 
dim3 numBlocks(2*N/threadsPerBlock.x); 
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,); 

(자원은 안, 허용 간단한 벡터 추가의 경우 문제가 될 수 있음). 따라서 더 나은 성능.

두 번째와 동일한 성능을 얻으려면 첫 번째 예제에서 블록 수를 20으로 곱하면되지만 두 번째 예제의 패턴은 사용 된 GPU의 특정 구성을 나타내지 않으므로 선호됩니다 . 따라서 코드는 넓은 범위의 GPU를 완전히로드합니다.

메모리 결합 알고리즘으로서의 벡터 추가는 점유율의 영향을 증명하는 데 특히 적합하지 않습니다. 그러나 메모리 하위 시스템 (메모리 대역폭과 메모리 액세스 대기 시간으로 결정)을 완전히로드하려면 비행 중 최소 최소 메모리 트랜잭션 수가 필요하며 5 % 점유 사례는이 최소값에 미치지 못하기 때문에 여전히 차이가 있습니다 .