내가 본 감소하는 동안

Question

대부분의 감소를 unactive 스레드를 종료 보면 같은 :

 
for(i = N; i > 0; i /= 2) { 
    if(tid >= i) 
     return; 
    assign-shared; 
    __syncthreads(); 
} 
copy-value-to-global; 

과 상당한 성능 이점을 발견 :

 
for(i = N; i > 0; i /=2) { 
    if(tid < i) 
     assign-shared; 
    __syncthreads(); 
} 
if(tid == 0) 
    copy-value-to-global; 

난 그냥 것을 반대했습니다. 축소가 더 이상 포함되지 않는 스레드를 갖는 것이 단점이 있습니까?

Answer 1

원래 코드로 if 문을 이미 수행 중이므로 어떤 단점도 보이지 않습니다.

if 구문의 결과에 공간적 지역성 (일반적으로 블록에서 동일한 결과)이없는 경우 속도가 향상되지 않을 수 있습니다. 또한 속도 향상은 장치 성능에 달려 있습니다. 이전 CUDA 장치는 성능 향상을 제공하지 않을 수 있습니다.

Answer 2

두 번째 코드 세그먼트는 사용되지 않은 워프가 다시 돌아와 분기 검사를 수행 할 필요가 없으므로 성능이 향상됩니다.

이상적으로, 두 번째 경우에는 반복 당 하나의 워프를 사용 중지하면 GPU의로드가 줄어 듭니다.

Answer 3

위의 그의 의견에서 dolan은 Can I use __syncthreads() after having dropped threads?에 따르면 윌리엄 퍼셀 (William Pursell)이 제안한 계획이 교착 상태에 빠질 것이라고 문제를 제기하고 있습니다. 이 문제에 관해서는 conditional syncthreads & deadlock (or not)에 따르면 코드가 대부분의 GPU에서 교착 상태가되지 않는다고합니다. 그 이유는 조기 종료를 지원하기 때문입니다. 하드웨어가 각 블록에 대한 활성 스레드 수를 유지하기 때문입니다.이 수는 장벽 동기화에 사용됩니다 블록의 초기 스레드 수보다

나는 reduce4 CUDA SDK 예제를 고려했으며 OP 질문에 따라 수정했습니다. 즉, 나는 두 __global__ 기능을 비교하고 있습니다 :

ORIGINAL

template <class T> 
__global__ void reduce4(T *g_idata, T *g_odata, unsigned int N) 
{ 
    extern __shared__ T sdata[]; 

    unsigned int tid = threadIdx.x;        // Local thread index 
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  // Global thread index - Fictitiously double the block dimension 

    // --- Performs the first level of reduction in registers when reading from global memory. 
    T mySum = (i < N) ? g_idata[i] : 0; 
    if (i + blockDim.x < N) mySum += g_idata[i+blockDim.x]; 
    sdata[tid] = mySum; 

    // --- Before going further, we have to make sure that all the shared memory loads have been completed 
    __syncthreads(); 

    // --- Reduction in shared memory. Only half of the threads contribute to reduction. 
    for (unsigned int s=blockDim.x/2; s>32; s>>=1) 
    { 
     if (tid < s) { sdata[tid] = mySum = mySum + sdata[tid + s]; } 
     // --- At the end of each iteration loop, we have to make sure that all memory operations have been completed 
     __syncthreads(); 
    } 

    // --- Single warp reduction by loop unrolling. Assuming blockDim.x >64 
    if (tid < 32) { 
     sdata[tid] = mySum = mySum + sdata[tid + 32]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 16]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 8]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 4]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 2]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 1]; __syncthreads(); 
    } 

    // --- Write result for this block to global memory. At the end of the kernel, global memory will contain the results for the summations of 
    //  individual blocks 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

이

template <class T> 
__global__ void reduce4_deadlock_test(T *g_idata, T *g_odata, unsigned int N) 
{ 
    extern __shared__ T sdata[]; 

    unsigned int tid = threadIdx.x;        // Local thread index 
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  // Global thread index - Fictitiously double the block dimension 

    // --- Performs the first level of reduction in registers when reading from global memory. 
    T mySum = (i < N) ? g_idata[i] : 0; 
    if (i + blockDim.x < N) mySum += g_idata[i+blockDim.x]; 
    sdata[tid] = mySum; 

    // --- Before going further, we have to make sure that all the shared memory loads have been completed 
    __syncthreads(); 

    // --- Reduction in shared memory. Only half of the threads contribute to reduction. 
    for (unsigned int s=blockDim.x/2; s>32; s>>=1) 
    { 
     if (tid >= s) return; 
     sdata[tid] = mySum = mySum + sdata[tid + s]; 
     // --- At the end of each iteration loop, we have to make sure that all memory operations have been completed 
     __syncthreads(); 
    } 

    // --- Single warp reduction by loop unrolling. Assuming blockDim.x >64 
    if (tid < 32) { 
     sdata[tid] = mySum = mySum + sdata[tid + 32]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 16]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 8]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 4]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 2]; __syncthreads(); 
     sdata[tid] = mySum = mySum + sdata[tid + 1]; __syncthreads(); 
    } 

    // --- Write result for this block to global memory. At the end of the kernel, global memory will contain the results for the summations of 
    //  individual blocks 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
    } 

수정 내가 수정 된 코드는 GT210, GT540M에 교착하지 않음을 확인하고있다 케플러 K20c. 그러나 케플러 카드, 수정 된 버전의 속도 향상은 관련 (ms에서 회)되지 않습니다 : 내가 다른 아키텍처에 대한 타이밍을 확인하지 않은

N   Original   Modified 
131072  0.021    0.019 
262144  0.030    0.032 
524288  0.052    0.052 
1048576 0.091    0.080 
2097152 0.165    0.146 
4194304 0.323    0.286 
8388608 0.637    0.555 
16777216 1.264    1.122 
33554432 2.514    2.189 

, 그러나 아마 위험은에 붙어 하락 일부 GPU의 교착 상태는 도달 가능한 속도 향상에 도달 할 가치가 없습니다 (도달 가능한 속도 향상이 동일한 순서로 유지되는 경우).