CUDA 버전이 CPU 버전보다 느립니다?

Question

CUDA에서 이미지 서브 샘플러를 작성하고 스레드를 사용하여 평균화 작업을 수행합니다. 그러나 커널을 호출하지 않고이 작업을 수행하면 실제로 CUDA 커널을 호출 할 때보 다 훨씬 빠르게 실행됩니다. 이미지 크기는 지금 1280x1024입니다. 커널 호출은 보통 상당한 시간이 걸리거나 구현에 문제가 있습니까?

피씨 나는 코드를 제거한 상태에서 커널을 호출 해 보았습니다. 코드가있는 커널과 거의 같은 시간이었습니다. 또한 커널 호출이 없으면 내 코드는 약 350ms를 실행하지만 커널 호출은 1000ms에 가깝습니다.

__global__ void subsampler(int *r_d,int *g_d,int *b_d, int height,int width,int *f_r,int*f_g,int*f_b){ 
     int id=blockIdx.x * blockDim.x*blockDim.y+ threadIdx.y*blockDim.x+threadIdx.x+blockIdx.y*gridDim.x*blockDim.x*blockDim.y; 
     if (id<height*width/4){ 
     f_r[id]=(r_d[4*id]+r_d[4*id+1]+r_d[4*id+2]+r_d[4*id+3])/4; 
     f_g[id]=(g_d[4*id]+g_d[4*id+1]+g_d[4*id+2]+g_d[4*id+3])/4; 
     f_b[id]=(b_d[4*id]+b_d[4*id+1]+b_d[4*id+2]+b_d[4*id+3])/4; 
     } 
     } 

나는 blockSizeX 및 blockSizeY 1 1 (내가 그들 4,16 만드는 시도)하지만, 어떻게 든이 가장 빠른

dim3 blockSize(blocksizeX,blocksizeY); 
    int new_width=img_width/2; 
    int new_height=img_height/2; 

    int n_blocks_x=new_width/blocksizeX+(new_width/blocksizeY == 0 ?0:1); 
    int n_blocks_y=new_height/blocksizeX+(new_height/blocksizeY == 0 ?0:1); 
    dim3 gridSize(n_blocks_x,n_blocks_y); 

하고 그때 gridSize, 블록 크기 커널 호출로 정의합니다.

Answer 1

커널이 제대로 구현되지 않았거나 GPU 카드에서 데이터를 이동하는 오버 헤드가 계산상의 이점을 저해 할 수 있습니다. (CPU < -> GPU 메모리 전송없이) 커널을 벤치마킹하여 전체 시간 중 얼마나 많은 부분을 커널이 차지하고 얼마나 많은 메모리를 전송했는지 확인하십시오. 커널에 대해 더 많은 작업을해야하는지 여부를 결정할 수 있습니다.

Answer 2

난 당신이 하나를 실행중인 하드웨어 확실하지 않다 있지만,이 커널 오히려 1000MS/프레임 :

제안 (1)보다 1000 FPS에 가깝게 수행 할 수 있어야한다 :이 처리하는 경우 OpenGL/DirectX 또는 이와 유사한 방법으로 시각화와 상호 작용합니다. 그리드/블록 크기, 메모리 레이아웃 등의 모든 세부 정보가 처리됩니다.

먼저 CUDA에서 1280x1024 이미지를 각 방향으로 2 배씩 서브 샘플링하여 640x512 이미지를 얻는다고 가정합니다. 결과 이미지의 각 픽셀은 원본 이미지의 4 픽셀 평균입니다. 이미지에는 RGB의 세 가지 채널이 있습니다.

질문 1 : 당신이 정말로 채널 당 32 비트 가 또는 RGB888 (채널 당 8 비트)를 원하는 않았다 하시겠습니까? RGB888은 매우 일반적입니다. 나는 이것이 당신이 의미하는 것이라고 생각합니다.

질문 2 : 데이터가 실제로 평면입니까, 아니면 인터리브 형식에서 추출하고 있습니까? RGB888은 픽셀이 RGBRGBRGB로 메모리에 나타나는 인터리브 형식입니다. 네이티브 형식으로 이미지를 처리하도록 커널을 작성합니다. 귀하의 데이터가 실제로 평면이라는 가정하에 R8, G8 및 B8 세 개의 평면이 있습니다.

먼저 메모리 레이아웃을 고려해야합니다. 대상 이미지의 모든 픽셀에 대해 하나의 스레드가 필요합니다. 서브 샘플링을위한 메모리 액세스 패턴이 통합되지 않는다면 픽셀 데이터를 공유 메모리로 읽어 들이기를 원할 것입니다. 32x8 스레드의 블록 크기를 고려하십시오. 이를 통해 각 블록은 40 * 8 * 4 픽셀 또는 3bpp에서 3072 바이트로 읽을 수 있습니다. 블록 당 총 4096 바이트의로드가 병합되도록 실제로 읽기보다 약간 더 많이 읽습니다. 이제 다음이 제공됩니다 :

dim3 block(32, 8); 
dim3 grid(1280/2/32, 1024/2/8); // 20x64 blocks of 256 threads 

이제 공유 메모리를 수행하는 재미있는 부분이 있습니다.커널은 다음과 같이 보일 수 있습니다 :

__global__ void subsample(uchar* r, uchar* g, uchar* b, // in 
          uchar* ro, uchar* go, uchar* bo) // out 
{ 
    /* Global offset into output pixel arrays */ 
    int gid = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x; 

    /* Global offset into input pixel arrays */ 
    int gidin = gid * 2; 

    __shared__ uchar* rc[1024]; 
    __shared__ uchar* gc[1024]; 
    __shared__ uchar* bc[1024]; 

    /* Read r, g, and b, into shmem cache */ 
    ((int*)rc)[threadIdx.x] = ((int*)r)[gidin + threadIdx.x]; 
    ((int*)gc)[threadIdx.x] = ((int*)g)[gidin + threadIdx.x]; 
    ((int*)bc)[threadIdx.x] = ((int*)b)[gidin + threadIdx.x]; 

    __syncthreads(); 

    /* Shared memory for output */ 
    __shared__ uchar* roc[256]; 
    __shared__ uchar* goc[256]; 
    __shared__ uchar* boc[256]; 

    /* Do the subsampling, one pixel per thread. Store into the output shared memory */ 

    ... 

    __syncthreads(); 

    /* Finally, write the result to global memory with coalesced stores */ 
    if (threadIdx.x < 64) { 
     ((int*)ro)[gid + threadIdx.x] = ((int*)roc)[threadIdx.x]; 
    } else if (threadIdx.x < 128) { 
     ((int*)go)[gid + threadIdx.x-64] = ((int*)goc)[threadIdx.x-64]; 
    } else if (threadIdx.x < 192) { 
     ((int*)bo)[gid + threadIdx.x-128] = ((int*)boc)[threadIdx.x-128]; 
    } 
} 

휴! 거기에 많은 것들이 있습니다. 코드 덤프는 유감스럽게 생각합니다. 명심해야 할 몇 가지 원칙 :

1) 병합 된로드/저장소를 사용할 때 메모리가 빠릅니다. 이는 32 개의 워프에있는 각 스레드가 각각 32 바이트를 액세스한다는 것을 의미합니다. 32 바이트 인덱스가 워프의 스레드 인덱스와 일치하면 32 개의 모든 액세스가 하나의 128 트랜잭션으로 처리됩니다. 이것은 GPU의 100GB/s 대역폭을 얻는 방법입니다.

2) 서브 샘플링을 수행 할 때 메모리 액세스 패턴은 병합되지 않습니다. 원시 메모리에없는 2D 공간 지역성에 의존하기 때문입니다. (이것에도 텍스처 메모리를 사용할 수 있습니다 ...) 입력을 공유 메모리에 저장 한 다음 처리함으로써 계산 성능에 미치는 영향을 최소화합니다.

이 정보가 도움이되기를 바랍니다. 원하실 경우 일부 부분에 대해 자세히 답변 해 드리겠습니다.