cudaMemcpy2D를 이용한 초간단 zero-padding

오랫만에 cuda 관련 포스트(지만 그닥 쓸모는 없어보인다.)

어떤 알고리즘을 구현하다보면 (특히 dsp알고리즘) 뒤에 0을 붙여주는 과정이 필요할 때가 생기는데, 이를 zero-padding이라 한다. 예를 들어 rgb데이터를 rgba데이터로 표시하고 싶다하면, 메모리상에는 linear로 rgbrgbrgbrgb... 이렇게 가지만, rgb0rgb0rgb0rgb0...이렇게 한 바이트씩 더 주고 싶은 경우라던가.

cpu에서 한다면 for문을 한개 돌려서 이런 식으로 할 수 있다.

mem=0, src=0;
for (int i=0; i< n; i++)
{
tempimg[mem++] = img1[src++];
tempimg[mem++] = img1[src++];
tempimg[mem++] = img1[src++];
tempimg[mem++] = 0; // padding zeros
}


cudaMemcpy2D를 이용하면 이것을 한 줄로 쓸 수 있다. 일단 cudaMemcpy2D 함수의 문법을 보면,

cudaError_t cudaMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);

여기서 spitch와 dpitch를 다르게 주는 것으로 가능. 위의 for문은 간단히 이렇게 쓸 수 있다.

cudaMemcpy2D(tempimg, 4, img1, 3, 3, n, cudaMemcpyHostToHost);


주의할 점은, 한 줄로 간단히 쓴다고 해서 더 빠르다는 것은 아니라는 점. 실제로 for문을 돌리는 것이 더 빨랐더랬다. 그치만 gpu메모리에 다시 memcpy를 해야하는 경우라면 좀 다를지도 모르겠다.

(cuda 포럼에 가보면 cudaMemcpy2D 함수가 느리다는 불평이 꽤나 많음.)