进出设备内存的结构深拷贝

Question

我在使用此 cuda 代码中动态分配的成员变量的结构数组的深层副本时遇到问题。我认为这是因为&deviceHistogram 指向主机上的地址而不是设备上的地址。我尝试在here 中创建一个中间指针变量，但这不起作用；如何正确复制整个结构数组，以便可以从 makeHistogram 函数对其进行修改？

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

typedef struct histogramBin {
    int* items;
    int count;
} histogramBin;

__host__ __device__ void outputHistogram(histogramBin* histogram, int size) {
    for (int i = 0; i < size; i++) {
        printf("%d: ", i);
        if (!histogram[i].count) {
            printf("EMPTY");
        } else {
            for (int j = 0; j < histogram[i].count; j++) {
                printf("%d ", histogram[i].items[j]);
            }
        }
        printf("\n");
    }
}


// This function embeds PTX code of CUDA to extract bit field from x. 
   __device__ uint bfe(uint x, uint start, uint nbits) {
    uint bits;
    asm("bfe.u32 %0, %1, %2, %3;"
        : "=r"(bits)
        : "r"(x), "r"(start), "r"(nbits));
    return bits;
}

__global__ void makeHistogram(histogramBin** histogram, int* rH, int rSize, int bit) {
    for (int r = 0; r < rSize; r++) {
        int thisBin = bfe(rH[r], bit, 1);
        int position = (*histogram)[thisBin].count; // **** out of memory access here****
        (*histogram)[thisBin].items[position] = rH[r];
        (*histogram)[thisBin].count++;
    }
}

void histogramDriver(histogramBin* histogram, int* rH, int rSize, int bit) {
    int n = 8;
    int* deviceRH;
    histogramBin* deviceHistogram;

    cudaMalloc((void**)&deviceRH, rSize * sizeof(int));
    cudaMemcpy(deviceRH, rH, rSize * sizeof(int), cudaMemcpyHostToDevice);

    cudaMalloc((void**)&deviceHistogram, n * sizeof(histogramBin));
    cudaMemcpy(deviceHistogram, histogram, n * sizeof(histogramBin), cudaMemcpyHostToDevice);

    int* tempData[n];
    for (int i = 0; i < n; i++) {
        cudaMalloc(&(tempData[i]), rSize * sizeof(int));
    }
    for (int i = 0; i < n; i++) {
        cudaMemcpy(&(deviceHistogram[i].items), &(tempData[i]), sizeof(int*), cudaMemcpyHostToDevice);
    }
    for (int i = 0; i < n; i++) {
        cudaMemcpy(tempData[i], histogram[i].items, rSize * sizeof(int), cudaMemcpyHostToDevice);
    }

    makeHistogram<<<1, 1>>>(&deviceHistogram, deviceRH, rSize, bit);
    cudaDeviceSynchronize();
}


int main(){
    int rSize = 5;
    int rH[rSize] = {1, 2, 3, 4, 5};

    histogramBin * histogram = (histogramBin*)malloc(sizeof(histogramBin) * 8);
    for(int i = 0; i < 8; i++){
        histogram[i].items = (int*)calloc(sizeof(int), rSize);
        histogram[i].count = 0;
    }
    histogramDriver(histogram, rH, rSize, 0);
    return 0;
}

一旦它被正确复制到设备上，我如何将它恢复到主机上？例如，如果我从 makeHistogram 内部调用 outputHistogram(histogram, 5);，我会看到以下内容：

0: 2 4 
1: 1 3 5 
2: EMPTY
3: EMPTY
4: EMPTY
5: EMPTY
6: EMPTY
7: EMPTY

这是我期望的输出。

当我从histogramDriver（cudaDeviceSynchronize() 之后）呼叫outputHistogram(histogram, 8) 时，我看到以下内容：

0: EMPTY
1: EMPTY
2: EMPTY
3: EMPTY
4: EMPTY
5: EMPTY
6: EMPTY
7: EMPTY

显然我没有正确地将值从设备复制回主机。

我尝试通过与histogramDriver 中的相反的过程进行复制：

for(int i = 0; i < n; i++){
    cudaMemcpy(&(tempData[i]), &(deviceHistogram[i].items), sizeof(int*), cudaMemcpyDeviceToHost);
}
for (int i = 0; i < n; i++) {
    cudaMemcpy(histogram[i].items, tempData[i], rSize * sizeof(int), cudaMemcpyDeviceToHost);
}

但histogramDriver 中的outputHistogram 调用的输出保持不变。

Answer 1

正如@talonmies 所指出的，这里最大的问题是内核的设计。没有理由/不需要为 histogram 使用双指针（事实上，您发布的代码的第一次迭代在内核原型中没有，尽管它不完整）。

通过删除双指针方面，您的代码运行时不会出现任何运行时错误。

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

typedef struct histogramBin {
    int* items;
    int count;
} histogramBin;

// This function embeds PTX code of CUDA to extract bit field from x.
   __device__ uint bfe(uint x, uint start, uint nbits) {
    uint bits;
    asm("bfe.u32 %0, %1, %2, %3;"
        : "=r"(bits)
        : "r"(x), "r"(start), "r"(nbits));
    return bits;
}

__global__ void makeHistogram(histogramBin* histogram, int* rH, int rSize, int bit) {
    for (int r = 0; r < rSize; r++) {
        int thisBin = bfe(rH[r], bit, 1);
        int position = histogram[thisBin].count; 
        histogram[thisBin].items[position] = rH[r];
        histogram[thisBin].count++;
    }
}

void histogramDriver(histogramBin* histogram, int* rH, int rSize, int bit) {
    int n = 8;
    int* deviceRH;
    histogramBin* deviceHistogram;

    cudaMalloc((void**)&deviceRH, rSize * sizeof(int));
    cudaMemcpy(deviceRH, rH, rSize * sizeof(int), cudaMemcpyHostToDevice);

    cudaMalloc((void**)&deviceHistogram, n * sizeof(histogramBin));
    cudaMemcpy(deviceHistogram, histogram, n * sizeof(histogramBin), cudaMemcpyHostToDevice);

    int* tempData[n];
    for (int i = 0; i < n; i++) {
        cudaMalloc(&(tempData[i]), rSize * sizeof(int));
    }
    for (int i = 0; i < n; i++) {
        cudaMemcpy(&(deviceHistogram[i].items), &(tempData[i]), sizeof(int*), cudaMemcpyHostToDevice);
    }
    for (int i = 0; i < n; i++) {
        cudaMemcpy(tempData[i], histogram[i].items, rSize * sizeof(int), cudaMemcpyHostToDevice);
    }

    makeHistogram<<<1, 1>>>(deviceHistogram, deviceRH, rSize, bit);
    cudaDeviceSynchronize();
}


int main(){
    const int rSize = 5;
    int rH[rSize] = {1, 2, 3, 4, 5};

    histogramBin * histogram = (histogramBin*)malloc(sizeof(histogramBin) * 8);
    for(int i = 0; i < 8; i++){
        histogram[i].items = (int*)calloc(sizeof(int), rSize);
        histogram[i].count = 0;
    }
    histogramDriver(histogram, rH, rSize, 0);
    return 0;
}
$ nvcc t1452.cu -o t1452
$ cuda-memcheck ./t1452
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$

请注意，这里唯一的更改是内核代码本身，加上删除内核调用时的 & 符号，另外我在 rSize 的定义中添加了 const 以进行编译。

我不知道它是否会产生正确的输出，因为您没有包含任何检查输出的方法，也没有指出您期望的输出是什么。如果您对此感兴趣，那么可以将这些内容包含在您的 MVE 中。