c＃managedCuda二維數組到GPU

Question

我是CUDA的新手，並試圖弄清楚如何將2d數組傳遞給內核。 我有以下工作代碼爲一維數組：

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 
    } 
} 

用下面的內核代碼：

__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; 
} 

正如我所說的，一切工作正常，但我想用二維數組如下工作：

// 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; 
    } 
} 

我需要的所有第二維是同一個線程，因此kernel.BlockDimensions必須保持爲1個捫sion和每個內核線程都需要獲得具有10個元素的1d數組。

所以我的底部問題是：我如何將這個2d數組複製到設備以及如何在內核中使用它？ （至於它應該總共有10個線程的例子）。

Answer 1

簡短的回答：你不應該這樣做...

龍答：鐵血陣列很難在一般的處理。你的數據中沒有連續的內存段，而是有很多小的內存在你內存中的某處。如果將數據複製到GPU會發生什麼？如果你有一個大的連續段，你可以調用cudaMemcpy/CopyToDevice函數並且一次複製整個塊。但是，與在for循環中分配鋸齒陣列時相同，您必須逐行將數據複製到CudaDeviceVariable<CUdeviceptr>，其中每個條目指向CudaDeviceVariable<float>。並行地維護主機端CudaDeviceVariable<float>[]，主機端管理你的CUdeviceptrs。複製數據一般已經很慢了，這樣做可能是一個真正的性能殺手...

總結：如果可以，請使用拼合數組並索引條目索引y * DimX + x。在GPU方面更好，使用傾斜內存，分配完成後，每行都在「好」地址處開始：然後索引變爲y * Pitch + x（簡化版）。 CUDA中的2D複製方法是爲這些傾斜的內存分配而製作的，其中每行添加一些額外的字節。

爲了完整性：在C＃中，您還有像2維數組，如float[,]。您也可以在主機端使用這些而不是扁平化的1D陣列。但我不建議這樣做，因爲.net的ISO標準不能保證內部存儲器實際上是連續的，ManageCuda必須使用這個假設才能使用這些陣列。目前.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]; 
}