在CUDA中分配设备内存上的2D数组

Question

如何在Cuda中分配和传输(到主机和从主机)设备内存中的2D数组？

- Gitmo

Answer 1

我找到了解决这个问题的方法。我不需要将数组压平。

内置的cudaMallocPitch()函数完成了这项工作。我可以使用cudaMemcpy2D()函数传输数组到设备和从设备传输数组。

例如：

cudaMallocPitch((void**) &array, &pitch, a*sizeof(float), b);

这将创建一个大小为a*b的2D数组，并将传递的参数作为间距。

以下代码创建一个2D数组并循环遍历元素。它可以轻松编译，您可以使用它。

#include<stdio.h>
#include<cuda.h>
#define height 50
#define width 50

// Device code
__global__ void kernel(float* devPtr, int pitch)
{
    for (int r = 0; r < height; ++r) {
        float* row = (float*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
             float element = row[c];
        }
    }
}

//Host Code
int main()
{    
    float* devPtr;
    size_t pitch;
    cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height);
    kernel<<<100, 512>>>(devPtr, pitch);
    return 0;
}

- Gitmo

---

是否可以稍后为数组分配新行？ - scatman

Answer 2

将其展平：使其变为一维数组。可以在这里查看如何操作。

- Rom

Answer 3

您的设备代码可以更快。尝试更充分地利用线程。

__global__ void kernel(float* devPtr, int pitch)
{
    int r = threadIdx.x;

    float* row = (float*)((char*)devPtr + r * pitch);
    for (int c = 0; c < width; ++c) {
         float element = row[c];
    }
}

然后，您需要计算适当的块和线程分配，以便每个线程处理一个单独元素。

- Abdullah

---

Gitmo发布的代码只是文档中无用的示例。是的，你的版本更快，但是如何并行处理行和列呢？严格来说，你可能会遇到麻烦，因为你没有检查r是否小于实际行数。 - darda

Answer 4

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h>
#include <cuda.h>
#define MAX_ITER 1000000
#define MAX 100 //maximum value of the matrix element
#define TOL 0.000001

// Generate a random float number with the maximum value of max
float rand_float(int max){
  return ((float)rand()/(float)(RAND_MAX)) * max;
}

__global__ void kernel(float **device_2Darray1, float **device_2Darray2, float **device_2Darray3, int rows, int cols) {

  // Calculate the row index
  int row = blockIdx.y * blockDim.y + threadIdx.y;

  // Calculate the column index
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  
  // Check if the thread is within the array bounds
  if (row < rows && col < cols) {
    // Perform the computation
    device_2Darray3[row][col] = device_2Darray1[row][col] + device_2Darray2[row][col];
  }
}


int main(int argc, char *argv[]){
  float **host_2Darray;
  float **device_2Darray;

  int rows = 10; // or whatever value you want
  int cols = 10; // or whatever value you want

  // allocate memory for the host
  host_2Darray = (float**)malloc(rows * sizeof(float*));
  for(int i = 0; i < rows; i++){
    host_2Darray[i] = (float*)malloc(cols * sizeof(float));
    for(int j = 0; j < cols; j++){
      host_2Darray[i][j] = rand_float(MAX);
    }
  }

  // allocate memory for the device
  cudaMalloc((void***)&device_2Darray, rows * sizeof(float*));
  for(int i = 0; i < rows; i++){
    cudaMalloc((void**)&device_2Darray[i], cols * sizeof(float));
  }

  // copy host memory to device
  for(int i = 0; i < rows; i++){
    cudaMemcpy(device_2Darray[i], host_2Darray[i], cols * sizeof(float), cudaMemcpyHostToDevice);
  }

  // call the kernel
  dim3 threadsPerBlock(16, 16);
  dim3 blocksPerGrid((rows + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                     (cols + threadsPerBlock.y - 1) / threadsPerBlock.y);
  kernel<<<blocksPerGrid, threadsPerBlock>>>(device_2Darray, rows, cols);

  // copy device memory back to host
  for(int i = 0; i < rows; i++){
    cudaMemcpy(host_2Darray[i], device_2Darray[i], cols * sizeof(float), cudaMemcpyDeviceToHost);
  }

  // free device memory
  for(int i = 0; i < rows; i++){
    cudaFree(device_2Darray[i]);
  }
  cudaFree(device_2Darray);

  // free host memory
  for(int i = 0; i < rows; i++){
    free(host_2Darray[i]);
  }
  free(host_2Darray);

  return 0;
}

在CUDA中创建2D数组的方法比1D数组更复杂，因为设备（GPU）内存是线性的。我们实际上创建了一个指针数组（每个指针指向一个1D数组），因此有双重指针。

至于全局void内核函数：像对两个2D数组逐元素相加这样的简单操作。您将使用blockIdx和blockDim以及threadIdx来计算线程的全局索引，然后使用此索引来指定该线程负责哪个数组元素。在此示例中，blockIdx.y和blockIdx.x给出了y和x方向上当前线程在块内的索引。您可以将其视为3级层次结构：网格->块->线程。

您需要调整代码的其余部分以分配和初始化第二个数组，分配输出数组，将输出复制回主机等。

不要忘记修改内核调用以传递正确的参数。您必须自己完成一些工作，例如集成概念，因为没有人可以访问您的完整代码以及它们如何作为系统一起运行。

- joe hoeller

---

我已经给你的答案授予了50个积分。然而，现在我发现源代码的某些部分丢失了。你测试过你的源代码吗？ - user366312

我更新了代码，很抱歉。 - joe hoeller