GPU 최적화 기초 - 분기문과 __syncthreads() 제거

알고리즘, 블록, 메모리, 분기문 중 분기문이다.

 

 

행렬과 벡터의 곱셈 후 블록 크기의 공유 메모리에 저장되어 있는 값을 합하는 과정에서 각각의 스레드 인덱스를 제어하여 필요한 값만을 합할 수 있다.

 

하지만, if문을 이용해 스레드를 제어하는 것 보다는

사용하게 되지 않는 값의 합까지 동일하게 병렬적으로 계산하도록 함으로 성능이 더 향상되어 if문을 제거하였다.

 

하프워프 단위 이하 스레드

공유메모리 접근시 읽기 및 쓰기 충돌로 인한 오류값이 발생하지 때문에, 동기화를 위한 __syncthreads() 함수 호출은 필요없다

 

 

그래서 알고리즘, 블록, 메모리, 분기문 최적화를 통해 무려 56배나 빨라진 계산속도를 볼 수 있다.

 

 

최적화된 행렬 곱셈 - 알고리즘, 블록, 메모리, 분기문 최적화

 

ThreadID = Tid.x;

blockID = Bdm.x;                                            // 블록, 스레드 1차원이니깨

__shared__ float temp[numThreadsPerBlock];   // 공유메모리 사용

 

// 각 스레드는 필요한 행의 계산을 한다

for (int y = blockID ; y < M의 세로행 크기 ; y+= 그리드 크기)   // 입력 행 세트

{                    

       for ( int x = threadID ; x < M의 가로열 크기 ; x+= 스레드 블록 크기)  // 입력 열 세트

            result += M의 (x, y)항 * V의 x항;           // 행렬 곱

       temp[threadID] = result;                           // 공유메모리에 M의 ID행, V의 ID열 결과저장

       __syncthreads();        // 스레드 동기화

 

       // 공유메모리에 저장된 값을 모음.

       temp[threadID] += temp[threadID + 8];

       temp[threadID] += temp[threadID + 4];

       temp[threadID] += temp[threadID + 2];

 

       if(32 이하의 threadID에 대해)

         temp[ 512 + (threadID>>4) ] = temp[threadID] + temp[threadID + 1];

         // 4비트 시프트하면 뭐가되냐, 앞이랑 뒤를 더하는데 아... 이해가 안된다.

       __syncthreads();

 

       if(16 이하의 threadID에 대해)

          temp[threadID] = temp[512 + threadID] + temp[528 + threadID];

          // __syncthreads()와 특정 threadID를 제한하는 if문이

          // halfwarp 이하의 단위에서 실행되기 때문에 필요하지 않다.

     

        temp[threadID] = temp[threadID] + temp[threadID + 8];

        temp[threadID] = temp[threadID] + temp[threadID + 4];

        temp[threadID] = temp[threadID] + temp[threadID + 2];

        temp[threadID] = temp[threadID] + temp[threadID + 1];

        W의 y항 = temp[0];         // 0번에 있는걸 넣네,, 이거 좀 구조가 다른가보다

     }

    __syncthreads();
}