CUDA - 2차원 배열의 할당과 이용

우선 잘못된 방법 중 하나로 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