logo标签列表

在GPU上为动态结构体数组进行内存分配

cstructcudadynamic-memory-allocation
3

我有一个问题，需要将结构体数组传递给GPU内核。我参考了这个主题 - cudaMemcpy分段错误，并编写了以下代码：

#include <stdio.h>
#include <stdlib.h>

struct Test {
    char *array;
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);
    for(int i = 0; i < n; i++)
        test[i].array = (char*)malloc(size * sizeof(char));

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);
    for(int i=0; i < n; i++) {
        cudaMalloc((void**)&(test[i].array), size * sizeof(char));
        cudaMemcpy(&(dev_test[i].array), &(test[i].array), size * sizeof(char), cudaMemcpyHostToDevice);
    }

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}

没有错误，但内核中显示的值是不正确的。我做错了什么？提前感谢任何帮助。

为什么是 cudaMalloc((void**)&(test[i].array), size * sizeof(char)); 而不是 cudaMalloc((void**)&(dev_test[i].array), size * sizeof(char));? 另外，应该使用 cudaMemcpy(dev_test[i].array, test[i].array, size * sizeof(char), cudaMemcpyHostToDevice);。 - francis
@Francis，它不起作用（分段错误（核心已转储））。在GPU上，我们无法以标准方式分配内存。 - Bakus123
1
附加友情建议：除非您理解了提问者面临的问题，否则不要从问题中提取代码...如果我的建议没有起作用，很抱歉。 我的建议是为 dev_test[i].array 分配内存，而不是为 test[i].array 分配内存；test[i].array 已经通过 test[i].array = (char*)malloc(size * sizeof(char)); 在 CPU 上分配了内存。 - francis
@francis，好的没问题。是的，test[i].array已经在CPU上分配了内存，但是在GPU上没有分配。我们无法为dev_test[i].array分配内存，因为这块内存只能从设备上访问。至少我是这样理解的。 - Bakus123
2个回答
11

  1. 这将为主机内存分配一个新指针：

 test[i].array = (char*)malloc(size * sizeof(char));

  • 这是将数据复制到主机内存中的那个区域：

    •  memcpy(test[i].array, temp, size * sizeof(char));

  • 这是用一个指向设备内存的新指针覆盖了之前分配给主机内存的指针（来自上面第1步），进行覆写操作：

    •  cudaMalloc((void**)&(test[i].array), size * sizeof(char));
    在第三步之后，你在第二步中设置的数据将完全丢失，并且无法以任何方式访问。参考你提供的问题/回答中的第三步和第四步：
    3. 在主机上创建一个名为myhostptr的独立int指针 4. 在设备上使用cudaMalloc为myhostptr分配int存储空间
    你没有这样做。你没有创建一个单独的指针。你重复使用了（擦除，覆盖）指向主机上你关心的数据的现有指针。这个问题/回答，也是你提供的答案中链接的，几乎完全给出了你需要按顺序遵循的步骤（在代码中）。
    下面是你的代码的修改版本，它正确地实现了你根据你提供的问题/回答中未正确实现的第三步、第四步和第五步：（参见标记第3、4、5步的注释）
    $ cat t755.cu
#include <stdio.h>
#include <stdlib.h>

struct Test {
    char *array;
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);
    for(int i = 0; i < n; i++)
        test[i].array = (char*)malloc(size * sizeof(char));

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

    // Step 3:
    char *temp_data[n];
    // Step 4:
    for (int i=0; i < n; i++)
      cudaMalloc(&(temp_data[i]), size*sizeof(char));
    // Step 5:
    for (int i=0; i < n; i++)
      cudaMemcpy(&(dev_test[i].array), &(temp_data[i]), sizeof(char *), cudaMemcpyHostToDevice);
    // now copy the embedded data:
    for (int i=0; i < n; i++)
      cudaMemcpy(temp_data[i], test[i].array, size*sizeof(char), cudaMemcpyHostToDevice);

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}

$ nvcc -o t755 t755.cu
$ cuda-memcheck ./t755
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$

    由于以上方法对于初学者来说可能会有一定的挑战性，通常的建议是不要使用它，而是 扁平化 数据结构。 扁平化通常意味着重新排列数据存储方式，以便删除必须单独分配的嵌套指针。

    扁平化此数据结构的一个微不足道的示例将使用此内容：

    struct Test {
    char array[5];
};

    当然，这种特定方法并不适用于每个目的，但它应该说明一般的想法/意图。通过这种修改，例如，代码变得更简单：

    $ cat t755.cu
#include <stdio.h>
#include <stdlib.h>

struct Test {
    char array[5];
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

int main(void) {

    int n = 4, size = 5;
    Test *dev_test, *test;

    test = (Test*)malloc(sizeof(Test)*n);

    for(int i=0; i < n; i++) {
        char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
        memcpy(test[i].array, temp, size * sizeof(char));
    }

    cudaMalloc((void**)&dev_test, n * sizeof(Test));
    cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

    kernel<<<1, 1>>>(dev_test);
    cudaDeviceSynchronize();

    //  memory free
    return 0;
}
$ nvcc -o t755 t755.cu
$ cuda-memcheck ./t755
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$
    谢谢。 “flatten your data structures” 是什么意思？ - Bakus123
    1
    我更新了我的答案以回应这个问题。但是，如果你在CUDA标签上搜索，你会找到许多关于“扁平化”的参考和示例。 - Robert Crovella
    1
    感谢@Robert Crovella的回答。上面的答案对我非常有用。我已经更新了代码，从使用字符数组改为使用整数数组结构体作为参考。在内核中，值被更新并返回到主机。
      #include <stdio.h>
  #include <stdlib.h>

  struct Test {
  int *array;
  };

  __global__ void kernel(Test *dev_test) {
  //    for(int i=0; i < 4; i++)
  //      for(int j=0; j < 5; j++) {
  //        printf("Kernel[X][i]: %d \n", dev_test[i].array[j]);
  //    }
  int i =threadIdx.x + blockIdx.x*blockDim.x;
  int j = threadIdx.y + blockIdx.y*blockDim.y;

  if(i<4 && j<5)
      dev_test[i].array[j] *= 2; 
  }

  int main(void) {

  int n = 4, size = 5;
  Test *dev_test, *test;

  test = (Test*)malloc(sizeof(Test)*n);
  for(int i = 0; i < n; i++)
      test[i].array = (int*)malloc(size * sizeof(int));

  for(int i=0;i<n;i++)
      for(int j=0; j<size; j++)
          test[i].array[j] = i*n+j;

  cudaMalloc((void**)&dev_test, n * sizeof(Test));
  cudaMemcpy(dev_test, test, n * sizeof(Test), cudaMemcpyHostToDevice);

  // Step 3:
  int *temp_data[n];
  // Step 4:
  for (int i=0; i < n; i++)
      cudaMalloc(&(temp_data[i]), size*sizeof(int));
  // Step 5:
  for (int i=0; i < n; i++)
      cudaMemcpy(&(dev_test[i].array), &(temp_data[i]), sizeof(int *), cudaMemcpyHostToDevice);
  // now copy the embedded data:
  for (int i=0; i < n; i++)
      cudaMemcpy(temp_data[i], test[i].array, size*sizeof(int), cudaMemcpyHostToDevice);

  dim3 threads(4,5);
  kernel<<<1, threads>>>(dev_test);
  cudaDeviceSynchronize();

  for(int i=0;i<n;i++)
      cudaMemcpy(test[i].array, temp_data[i], size*sizeof(int), cudaMemcpyDeviceToHost);

  for(int i=0;i <n;i++){
      for(int j=0;j<size;j++)
          printf("%d ",test[i].array[j]);
  printf("\n");
  }
  //  memory free
  return 0;
  }
    网页内容由stack overflow 提供, 点击上面的
    可以查看英文原文，
    原文链接