C＃managedCuda 2dアレイからGPUへ

Question

私はCUDAが新しく、2d配列をカーネルに渡す方法を理解しようとしています。 私は1次元配列のために、次の作業のコードを持っている：私はすべてが正常に動作しますが、私は次のように2次元配列で仕事をしたい、言ったように

__global__ void DoubleIt(const float* A, float* C, int N) 
{ 
    int i = blockDim.x * blockIdx.x + threadIdx.x; 
    if (i < N) 
     C[i] = A[i] * 2; 
} 

次のカーネルコードと

class Program 
{ 
    static void Main(string[] args) 
    { 
     int N = 10; 
     int deviceID = 0; 
     CudaContext ctx = new CudaContext(deviceID); 
     CudaKernel kernel = ctx.LoadKernel(@"doubleIt.ptx", "DoubleIt"); 
     kernel.GridDimensions = (N + 255)/256; 
     kernel.BlockDimensions = Math.Min(N,256); 

     // Allocate input vectors h_A in host memory 
     float[] h_A = new float[N]; 

     // Initialize input vectors h_A 
     for (int i = 0; i < N; i++) 
     { 
      h_A[i] = i; 
     } 

     // Allocate vectors in device memory and copy vectors from host memory to device memory 
     CudaDeviceVariable<float> d_A = h_A; 
     CudaDeviceVariable<float> d_C = new CudaDeviceVariable<float>(N); 

     // Invoke kernel 
     kernel.Run(d_A.DevicePointer, d_C.DevicePointer, N); 

     // Copy result from device memory to host memory 
     float[] h_C = d_C; 
     // h_C contains the result in host memory 
    } 
} 

私は kernel.BlockDimensionsが1 DIMENとして滞在しなければならないので、同じスレッド上にあるように、すべての第二の寸法を必要とする

// Allocate input vectors h_A in host memory 
int W = 10; 
float[][] h_A = new float[N][]; 

// Initialize input vectors h_A 
for (int i = 0; i < N; i++) 
{ 
    h_A[i] = new float[W]; 
    for (int j = 0; j < W; j++) 
    { 
     h_A[i][j] = i*W+j; 
    } 
} 

各カーネルスレッドは10要素で1d配列を取得する必要があります。

私の下の質問は次のようなものです：この2次元配列をデバイスにコピーしてカーネルでどのように使用するのですか？ （この例では合計10個のスレッドが必要です）。

Answer 1

短い答え：あなたがそれを行うべきではありません...

長い答え：ギザギザの配列は、一般的に取り扱いが困難です。あなたのデータのための1つの連続したメモリセグメントの代わりに、あなたは小さなものがまばらにあなたの記憶のどこかに横たわっています。 GPUにデータをコピーするとどうなりますか？大きな連続したセグメントがある場合は、cudaMemcpy/CopyToDevice関数を呼び出して、ブロック全体を一度にコピーします。しかし、forループでギザギザの配列を割り当てるのと同じように、データを行ごとにCudaDeviceVariable<CUdeviceptr>にコピーしなければなりません。各エントリはCudaDeviceVariable<float>を指します。並行して、ホスト側でCUdeviceptrsを管理するホストアレイCudaDeviceVariable<float>[]を維持します。データを一般的にコピーすることは既に非常に遅いので、このようにすれば実際のパフォーマンスが低下するでしょう。

結論：平坦化された配列を使用し、インデックスy * DimX + xでエントリをインデックス化します。 GPU側でも、ピッチメモリを使用してください。割り当てが行われると、各行は「良い」アドレスで開始されます。インデックスはy * Pitch + x（簡略化）に変わります。 CUDAの2Dコピーメソッドは、各行にいくつかの追加バイトが追加されたこれらのピッチメモリ割り当て用に作られています。

完全性：C＃では、float[,]のような2次元配列もあります。また、1D配列のフラット化ではなく、ホスト側でも使用できます。しかし、私はそうすることをお勧めしません.NetのISO標準は、内部メモリが実際に連続であることを保証していないため、managedCudaがこれらの配列を使用するために使用しなければならないという前提です。現在.NET Frameworkは、任意の内部すごみを持っていないが、それはこのように滞在する場合、誰が知っている...

これは、ジャグ配列のコピーを実現します：

float[][] data_h; 
CudaDeviceVariable<CUdeviceptr> data_d; 
CUdeviceptr[] ptrsToData_h; //represents data_d on host side 
CudaDeviceVariable<float>[] arrayOfarray_d; //Array of CudaDeviceVariables to manage memory, source for pointers in ptrsToData_h. 

int sizeX = 512; 
int sizeY = 256; 

data_h = new float[sizeX][]; 
arrayOfarray_d = new CudaDeviceVariable<float>[sizeX]; 
data_d = new CudaDeviceVariable<CUdeviceptr>(sizeX); 
ptrsToData_h = new CUdeviceptr[sizeX]; 
for (int x = 0; x < sizeX; x++) 
{ 
    data_h[x] = new float[sizeY]; 
    arrayOfarray_d[x] = new CudaDeviceVariable<float>(sizeY); 
    ptrsToData_h[x] = arrayOfarray_d[x].DevicePointer; 
    //ToDo: init data on host... 
} 
//Copy the pointers once: 
data_d.CopyToDevice(ptrsToData_h); 

//Copy data: 
for (int x = 0; x < sizeX; x++) 
{ 
    arrayOfarray_d[x].CopyToDevice(data_h[x]); 
} 

//Call a kernel: 
kernel.Run(data_d.DevicePointer /*, other parameters*/); 

//kernel in *cu file: 
//__global__ void kernel(float** data_d, ...) 

これはCudaPitchedDeviceVariableのサンプルです：

int dimX = 512; 
int dimY = 512; 
float[] array_host = new float[dimX * dimY]; 
CudaPitchedDeviceVariable<float> arrayPitched_d = new CudaPitchedDeviceVariable<float>(dimX, dimY); 
for (int y = 0; y < dimY; y++) 
{ 
    for (int x = 0; x < dimX; x++) 
    { 
     array_host[y * dimX + x] = x * y; 
    } 
} 

arrayPitched_d.CopyToDevice(array_host); 
kernel.Run(arrayPitched_d.DevicePointer, arrayPitched_d.Pitch, dimX, dimY); 

//Correspondend kernel: 
extern "C" 
__global__ void kernel(float* data, size_t pitch, int dimX, int dimY) 
{ 
    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 
    if (x >= dimX || y >= dimY) 
     return; 

    //pointer arithmetic: add y*pitch to char* pointer as pitch is given in bytes, 
    //which gives the start of line y. Convert to float* and add x, to get the 
    //value at entry x of line y: 
    float value = *(((float*)((char*)data + y * pitch)) + x); 

    *(((float*)((char*)data + y * pitch)) + x) = value + 1; 

    //Or simpler if you don't like pointers: 
    float* line = (float*)((char*)data + y * pitch); 
    float value2 = line[x]; 
}