2 행렬의 내적에 대한 감소 - 논리 오류가있는 Cuda 커널

Question

저는 CUDA로 시작하여 CUDA 감소 알고리즘에 대해 두뇌를 감싸하려고합니다. 제 경우에는 두 행렬의 내적을 구하려고 노력했습니다. 하지만 크기가 2 인 행렬에 대해서만 올바른 답을 얻고 있습니다. 다른 크기 행렬에 대해서는 잘못 이해하고 있습니다.

이것은 테스트이므로 매트릭스 크기를 매우 작게 유지합니다. 약 100 개 밖에 없으므로 1 블록만으로도 충분합니다. 도움을 주시면 대단히 감사하겠습니다. 감사! 여기

내 CUDA 커널 일반 코드 여기

float* ha = new float[n]; // matrix a 
float* hb = new float[n]; // matrix b 
float* hc = new float[1]; // sum of a.b 

float dx = hc[0]; 
float hx = 0; 
// dot product 
for (int i = 0; i < n; i++) 
    hx += ha[i] * hb[i]; 

입니다

__global__ void sum_reduce(float* da, float* db, float* dc, int n) 
{ 
    int tid = threadIdx.x; 
    dc[tid] = 0; 
    for (int stride = 1; stride < n; stride *= 2) { 
     if (tid % (2 * stride) == 0) 
       dc[tid] += (da[tid] * db[tid]) + (da[tid+stride] * db[tid+stride]); 
     __syncthreads(); 
    } 
} 

내 전체 코드 :가 N = 2 일 왜 http://pastebin.com/zS85URX5

Answer 1

바라건대 당신은 알아낼 수 이 경우를 건너 뛰고 다른 경우에 실패하는 이유를 살펴보고 n = 4를 선택하겠습니다. 당신은 4 개 스레드가있을 때 N = 4, 첫 번째에 대한 루프, 스트라이드가 1, 그래서 만약 시험을 통과 스레드 스레드 0입니다 당신의 반복과 2

에서 3

에 번호가 0

thread 0: dc[0] += da[0]*db[0] + da[1]*db[1]; 
thread 2: dc[2] += da[2]*db[2] + da[3]*db[3]; 

지금까지 그렇게 좋았습니다. for 루프의 두 번째 반복에서는 stride가 2이므로 if 테스트를 통과 한 스레드는 스레드 0 (유일한)이됩니다.

thread 0: dc[0] += da[0]*db[0] + da[2]*db[2]; 

그러나 이것은 의미가 없으며 우리가 원하는 것이 아닙니다. 우리가 원하는 것은 다음과 같습니다 :

dc[0] += dc[2]; 

그래서 망가졌습니다. 약간의 단계만으로이 문제를 해결할 수있는 방법을 생각하면서 잠시 시간을 보냈지 만, 감수성으로 이해가되지 않습니다. 커널 코드를이 코드로 대체하면 좋은 결과를 얻을 수 있다고 생각합니다. 그것은 당신의 코드와 많이 다르지는 않지만, 당신이 상상 한 모든 경우에 사용할 수있는 가장 가까운 것이 었습니다 (즉, 하나의 블록을 사용하여 최대 스레드 블록 크기 인 <).

// CUDA kernel code 
__global__ void sum_reduce(float* da, float* db, float* dc, int n) 
{ 
    int tid = threadIdx.x; 
    // do multiplication in parallel for full width of threads 
    dc[tid] = da[tid] * db[tid]; 
    // wait for all threads to complete multiply step 
    __syncthreads(); 
    int stride = blockDim.x; 
    while (stride > 1){ 
     // handle odd step 
     if ((stride & 1) && (tid == 0)) dc[0] += dc[stride - 1]; 
     // successively divide problem by 2 
     stride >>= 1; 
     // add each upper half element to each lower half element 
     if (tid < stride) dc[tid] += dc[tid + stride]; 
     // wait for all threads to complete add step 
     __syncthreads(); 
     } 
} 

나는 실제로 n 매개 변수를 사용하지 않습니다. n 쓰레드로 커널을 실행하기 때문에,이 경우에는 내장 변수가 n과 같습니다.