CUDA의 성능

Question

내 커널 성능에 영향을 미치는 근본적인 요인을 파악할 수 없습니다. 두 개의 이미지를로드하고 두 이미지를로드하는 두 개의 간단한 커널과 두 이미지를로드하는 두 개의 간단한 커널을 구현했으며 비트 감에서 ANDS을 추가했습니다. 이제 커널을 템플릿 화하여 커널이 8 비트 및 32 비트 이미지는 물론 1, 3, 4 채널 이미지를 사용할 수있게했습니다.

그래서 처음에 나는 모두 커널 uchar4 등 내가 때문에 병합의, 그러나, 트리플 사용에 대한 너무 확신하지 못했습니다과 함께 uchar3 및 float3으로 전역 메모리로드를했다, 그래서 나는 그것을 줄 거라고 생각 프로파일 링 실행. 나는 작업이 채널 번호와 독립적 이었기 때문에 실제로는 uchar3 이미지가 아니라 1 채널 uchar 이미지가 3 배 넓이 인 것처럼 이미지를 읽을 수 있다고 생각했습니다.

실제로 uchar3 전역로드는 로드보다 훨씬 느립니다. 나의 노력은 옳았다. 그러나 슬프게도 이것은 산술 커널에서만 발생했습니다. 비트 AND 연산은 정확한 결과를 보여줍니다!

자, 이제 uchar이 아닌 uint s 이미지 데이터를로드 할 수 있다는 것을 알고 있습니다. 비트 작업의 경우 완벽하게 병합 처리해야합니다. 그러나 나는 단순히 무슨 일이 일어나고 있는지 배우고 이해하기를 원한다고 가정 해 봅시다.

그리고 float3 및 float4 등을 잊어 보겠습니다. 내 문제는 커널의 uchar 버전입니다. 따라서 간단히 말해 uchar로드가 uchar3로드보다 더 빠르며 때로는로드되지 않는 경우가 있습니까?

저는 GTX 470의 컴퓨팅 기능 2.0을 사용하고 있습니다.

추신. CUDA 프로그래밍 가이드에 따르면 논리 연산과 추가 연산은 동일한 처리량을가집니다. (제 커널은 실제로 uchar을 uint으로 변환해야만합니다.하지만 두 커널 모두에서 커널이 발생해야합니다.) 그래서 실행 길이는 내가 모은 것과 거의 같아야합니다.

산술 (uchar 버전) 커널을 추가

__global__ void add_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep) 
{ 
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x; 
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y; 

    if (yCoordinate >= height) 
     return; 

#pragma unroll IMAGE_MULTIPLIER 
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i) 
    { 
     // Load memory. 
     uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 
     uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 

     // Write output. 
     *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] + inputElementTwo[0]; 
    } 
} 

비트 단위 AND 커널 :

__global__ void and_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep) 
{ 
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x; 
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y; 

    if (yCoordinate >= height) 
     return; 

#pragma unroll IMAGE_MULTIPLIER 
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i) 
    { 
     // Load memory. 
     uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 
     uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 

     // Write output. 
     *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] & inputElementTwo[0]; 
    } 
} 

uchar3 버전

다음과 같이로드/저장 라인이 이제 것을 제외하고는 동일하다 :

 // Load memory. 
    uchar3 inputElementOne = *reinterpret_cast<uchar3*>(inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3); 
    uchar3 inputElementTwo = *reinterpret_cast<uchar3*>(inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3); 

    // Write output. 
    *reinterpret_cast<uchar3*>(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3) 
     = make_uchar3(inputElementOne.x + inputElementTwo.x, inputElementOne.y + inputElementTwo.y, inputElementOne.z + inputElementTwo.z); 

마찬가지로 AND 커널. (나는 커널을 정확히 기억하는지 모르지만, 솔직히 말해서 나는 내일 그것을 확증 할 것이다.) SM에의 명령어 세트에는 24 비트로드가 없기 때문에

Answer 1

uchar3 하중은 별도의 하중에 컴파일러에 의해 분할됩니다. 따라서, 그들은 결코 합체되지 않습니다. 캐시가이를 어느 정도 완화시킬 것입니다.

정확한 실행 구성에 따라 스레드 당 약 10.7 바이트의 캐시가있을 수 있습니다. (예는 커널이 간단하기 때문에이 값은 그 값에 가까울 것입니다. 따라서 많은 스레드가 하나의 SM에서 동시에 실행될 수 있습니다).캐시가 완전히 연관되어 있지 않으므로 스 래싱이 발생하기 전에 스레드 당 사용 가능한 바이트 수가 훨씬 적을 수 있습니다. 정확히 일어날 때 정확히 같은 처리량을 가진 지침에 대해 다를 수도 명령의 정확한 일정을 포함하여 많은 요인에 따라 달라집니다.

실행 파일 cuobjdump -sass을 두 버전 모두 비교하여 컴파일러의 정적 스케줄링이 동일한 지 여부를 확인할 수 있습니다. 그러나 런타임에서의 동적 스케쥴링은 기본적으로 관찰 할 수 없습니다.

이미 알고 있듯이 이미지의 모든 채널은 동일한 방식으로 처리되므로 스레드간에 배포하는 방법은 중요하지 않습니다. 가장 좋은 방법은 uchar3 또는 uchar 대신 uchar4을 사용하는 것입니다. 이미지의 적절한 정렬로 인해 캐시와 독립적으로 액세스가 통합됩니다. 이렇게하면 실행 시간이 더 짧아지고 일관성이 유지됩니다.