CUDA의 멀티 스레드 커널에서 불필요한 쓰기 작업이 비효율적입니까?

Question

나는 내 CUDA 코드에 커널을 가지고 있는데, 공유 메모리의 일부분에 많은 쓰레드를 쓰고 싶다. (전역 메모리를 사용하는 것보다 훨씬 빠르기 때문이다.) 그 결과를 전역 메모리에 쓴다. (그래서 나는 나중에 커널에서 사용할 수있다).

__global__ void calc(float * globalmem) 
{ 
    __shared__ float sharemem; //initialize shared memory 
    sharemem = 0; //set it to initial value 
    __syncthreads(); 

    //do various calculations on the shared memory 
    //for example I use atomicAdd() to add each thread's 
    //result to sharedmem... 

    __syncthreads(); 
    *globalmem = sharedmem;//write shared memory to global memory 
} 

글로벌 메모리 공유으로부터 정말 한 번만를 작성해야 할 때마다 하나의 스레드가 데이터를 쓰고 있다는 사실, 나에게 비린내 느낌 : 커널은 다음과 같이 보입니다. 또한 모든 스레드가 공유 메모리를 코드 시작 부분에서 0으로 초기화한다는 사실에서 같은 느낌을받습니다. 현재 구현보다 빠른 방법이 있습니까?

Answer 1

워프 레벨에서 중복 읽기 또는 쓰기 대 단일 스레드를 갖는 것 사이의 성능 차이는 거의 없습니다.

그러나 스레드 블록에 여러 개의 워프가있을 경우 중복 읽기 또는 쓰기 (단일 스레드 대)로 인해 측정 가능한 성능 차이가 발생할 것으로 예상됩니다.

오히려 중복보다 단일 스레드를함으로써 이러한 문제를 해결 읽기을하거나 쓰기에 충분해야한다 : 공유에 당신이 threadblock 내 아토을 사용하여, 그것에 대해 물어 보지 않았지만

__global__ void calc(float * globalmem) 
{ 
    __shared__ float sharemem; //initialize shared memory 
    if (!threadIdx.x) sharemem = 0; //set it to initial value 
    __syncthreads(); 

    //do various calculations on the shared memory 
    //for example I use atomicAdd() to add each thread's 
    //result to sharedmem... 

    __syncthreads(); 
    if (!threadIdx.x) *globalmem = sharemem;//write shared memory to global memory 
} 

메모리는 아마도 공유 메모리 감소 방법에 의해 (더 나은 성능을 위해) 대체 가능할 수 있습니다.