3D Client Programmer

CUDA는 계산 능력이 1.1 이상의 버전에서 원자적 연산(automic operation)을 지원한다.

하지만, 원자적 연산을 모든 블록의 쓰레드가 동시에 실행하게 되면, 그 원자에 접근하려는 수천 개의 스레드들이 치열한 다툼을 하게 되면서 오히려 성능 저하가 발생하게 된다.

그럴 때에는 공유 메모리를 사용하여, 원자에는 최대한 접근을 하지 않도록 하는 것이 좋다.

NVIDIA에서 공식적으로 제공하는 방법이다.

위 코드와 같은 방법으로 사용이 가능하다.

우선 두개의 cudaEvent 변수를 생성해야 하는데, 이것은 record 하는 순간의 timestamp를 저장하는 형식이기 때문이다.

이후 cudaEventRecord()를 이용하여 시작하는 순간과 끝나는 순간의 timestamp를 저장하면 된다.

시간은 cudaEventElapsedTime()를 이용하여 받아오는데, cudaEvent는 기본적으로 float으로 반환한다.

cudaEventSynchronize()는 Host에서 사용되는 함수로써, event가 발생할 때까지 Host는 대기하게 된다.

이와 비슷한 함수로 cudaStreamWaitEvent()가 있는데, 이 함수는 Device내부에서 event가 발생할 때까지 대기한다.

우선 잘못된 방법 중 하나로 malloc()라서 간과하고 일반적인 malloc()로 2차원 배열을 선언하는 방식은 접근 에러가 뜨게 된다.

그 원인이 한번 선언된 순간 해당 메모리 주소 정보는 메인에서 관리되지 않는다고 한다. 즉 Host에서 익숙하게 malloc() 사용하듯이 cudaMalloc()을 사용하는 순간 Host(CPU 이하 Host)에 등록된 주소는 잃게되고 Device(GPU 이하 Device)의 주소로 넘어가게 되는데 이때 Host에서는 접근이 불가하게 된다.

근데 Host에서 명령하는 cudaMalloc()을 사용하게 되면 접근이 불가능한데 d_array의 주소를 Host가 접근을 하지만 이미 접근권한을 잃은 Host는 d_array에 대해 nullptr 혹은 잘못된 값을 받아오게 될 것이다. 잘못된 주소를 Device에게 메모리 할당을 요청하기 때문에 에러가 발생하게 된다.

마찬가지로 할당된 주소를 달아주는 순간 Host는 그  권한을 잃게된다...

해결방안1. (반만 해결됨)

1차원 배열로 선언하되 2차원 배열로 이용한다. 뭔소리인가 하면, x, y인 2d 데이터에 접근할 때 data[y * width + x]의 형태로 접근하는 것이다.

쉽게 말해서 2 X 5 행렬을 사용한다면 10칸을 할당해서 아래와 같이 논리적으로 생각하고 쓴다는 것이다.

해결방안2. (추천 방법)

Example:

// kernel which copies data from d_array to destinationArray
__global__ void CopyData(float* d_array, 
                                   float* destinationArray, 
                                   size_t pitch, 
                                   int columnCount, 
                                   int rowCount)
{
  for (int row = 0; row < rowCount; row++) 
  {
     // update the pointer to point to the beginning of the next row
     float* rowData = (float*)(((char*)d_array) + (row * pitch));
        
    for (int column = 0; column < columnCount; column++) 
    {
      rowData[column] = 123.0; // make every value in the array 123.0
      destinationArray[(row*columnCount) + column] = rowData[column];
    }
  }
}


int main(int argc, char** argv) 
{	
  int columnCount = 15; 
  int rowCount = 10;
  float* d_array; // the device array which memory will be allocated to
  float* d_destinationArray; // the device array
  
  // allocate memory on the host
  float* h_array = new float[columnCount*rowCount];

  // the pitch value assigned by cudaMallocPitch
  // (which ensures correct data structure alignment)
  size_t pitch; 
  
  //allocated the device memory for source array
  cudaMallocPitch(&d_array, &pitch, columnCount * sizeof(float), rowCount);
  
  //allocate the device memory for destination array
  cudaMalloc(&d_destinationArray,columnCount*rowCount*sizeof(float));
  
  //call the kernel which copies values from d_array to d_destinationArray
  CopyData<<<100, 512>>>(d_array, d_destinationArray, pitch, columnCount, rowCount);

  //copy the data back to the host memory
  cudaMemcpy(h_array,
                    d_destinationArray,
                    columnCount*rowCount*sizeof(float),
                    cudaMemcpyDeviceToHost);

  //print out the values (all the values are 123.0)
  for(int i = 0 ; i < rowCount ; i++)
  {
    for(int j = 0 ; j < columnCount ; j++)
    {
      cout << "h_array[" << (i*columnCount) + j << "]=" << h_array[(i*columnCount) + j] << endl;
    }
  }
}

참고 사이트 : http://www.stevenmarkford.com/allocating-2d-arrays-in-cuda/

출처 : http://k1321a.blog.me/220357402521

CUDA 코드는 앞에서 다 봤기 때문에 더 설명할 내용은 없는 듯 하다.

- 소스 코드 -

main.cpp

texture_2d.cu

이번엔 RunKernels()에 대해 한번 보자.

앞에서 cudaGraphicsMapResource()로 주소 공간을 할당한 것을 cudaGraphicsSubResourceGetMappedArray()를 이용하여 배열로 반환 받는다.

cuda_texture_2d()는 예제에서 구현된 함수로 다음 장에 작성하겠다.

cudaMemcpy2DToArray()는 이름대로 2차원 배열을 복사한다.

이 함수를 호출 함으로써 DirectX와 CUDA 양쪽에서 ID3D11Resource를 사용할 것이라고 CUDA 런타임에게 명시할 수 있다. CUDA 런타임은 변수 resource를 통해 버퍼를 가리키는 CUDA 전용 핸들을 하나 반환한다. 이 핸들은 차후에 CUDA 런타임의 API를 호출시 ID3D11Resource를 참조하기 위해 사용될 것이다.

이 함수는 아래 사이트에서 확인하자.

https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__D3D11_gf0545f2dd459ba49cfd6bcf2741c5ebd.html

cudaMallocPitch()는 CUDA에서 2차원 배열을 생성하는 함수이다. C처럼 2차원 배열을 만들 수 없기 때문에 이렇게 따로 함수가 있다. 이것에 대해서는 다음에 다시 글을 작성하도록 하겠다.

cudaMemset()은 말 그대로 값 세팅을 한다.

cudaGraphicsMapResources()는 앞에서 ppResource를 CUDA에 주소를 할당하고, cudaGraphicsUnmapResources()은 ppResource의 주소 할당을 해제한다.

앞의 함수들을 호출했다면, 이제 DirectX를 초기화 해야한다.

앞에서 g_pCudaCapableAdapter를 구했었다. 이 변수를 이용해서 D3D11의 Device를 생성해야한다.

이때 주의해야 할 점은 D3D11CreateDevice()를 사용할 때 첫 번째 인자로 Adapter를 지정했다면, 두 번째 인자는 D3D_DRIVER_TYPE_UNKNOWN으로 반드시 지정해야 한다.

그리고 CUDA Adapter를 이용하여 Device를 생성했다면, 이제 CUDA에도 D3D Device를 지정해줘야 한다.

이 함수는 ::IDXGIFactory::EnumAdapters에서 찾은 Adapter에 해당하는 CUDA 호환 장치를 *device에 넣어 반환한다.

이 함수는 Adapter가 CUDA를 지원해야만 성공적으로 cudaSuccess 리턴을 한다.

다음 함수는 Cuda를 지원하는 Device를 찾아 IDXGIAdapter의 변수 g_pCudaCapableAdapter에 저장하는 함수이다.

이 g_pCudaCapableAdapter는 이후 D3D11CreateDeviceAndSwapChain()를 할 때 사용된다,