CUDA:将统一内存与类和数组一起使用

CUDA: using unified memory together with classes and arrays

本文关键字：数组 一起 内存 CUDA      更新时间：2023-10-16

我正在尝试获得统一内存以处理类，并通过内核调用在统一内存中传递和操作数组。我想通过引用传递所有内容。

所以我正在覆盖类和数组的新方法，以便 GPU 可以访问它们，但我认为我需要添加更多代码才能在统一内存中拥有数组，但不太确定如何做到这一点。调用 fillArray（） 方法时出现内存访问错误。

如果我必须执行数百次此类操作（数组上的算术和在不同大小的数组之间进行复制），统一内存是一种好方法还是应该坚持在 CPU 和 GPU 内存之间手动复制？谢谢！

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <stdio.h>

#define TILE_WIDTH 4
#ifdef __CUDACC__
#define CUDA_CALLABLE_MEMBER __host__ __device__
#else
#define CUDA_CALLABLE_MEMBER
#endif
__global__ void add1(int height, int width, int *a, int *resultArray)
{
    int w = blockIdx.x * blockDim.x + threadIdx.x; // Col // width
    int h = blockIdx.y * blockDim.y + threadIdx.y;
    int index = h * width + w;
    if ((w < width) && (h < height))
        resultArray[index] = a[index] + 1;
}
class Managed 
{
public:
    void *operator new(size_t len) 
    {
        void *ptr;
        cudaMallocManaged(&ptr, len);
        return ptr;
    }
    void Managed::operator delete(void *ptr) 
    {
        cudaFree(ptr);
    }
    void* operator new[] (size_t len) {
        void *ptr; 
        cudaMallocManaged(&ptr, len);
        return ptr;
    }
        void Managed::operator delete[] (void* ptr) {
        cudaFree(ptr);
    }
};
class testArray : public Managed
{
public: 
    testArray()
    {
        height = 16;
        width = 8;
        myArray = new int[height*width];
    }
    ~testArray()
    {
        delete[] myArray;
    }
    CUDA_CALLABLE_MEMBER void runTest()
    {
        fillArray(myArray);
        printArray(myArray);
        dim3 dimGridWidth((width - 1) / TILE_WIDTH + 1, (height - 1)/TILE_WIDTH + 1, 1);
        dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
        add1<<<dimGridWidth,dimBlock>>>(height, width, myArray, myArray);
        cudaDeviceSynchronize();
        printArray(myArray);
    }
private:
    int *myArray;
    int height; 
    int width;
    void fillArray(int *myArray)
    {
        for (int i = 0; i < height; i++){
            for (int j = 0; j < width; j++)
                myArray[i*width+j] = i*width+j;
        }
    }
    void printArray(int *myArray)
    {
        for (int i = 0; i < height; i++){
            for (int j = 0; j < width; j++)
                printf("%i ",myArray[i*width+j]);
            printf("n");
        }
    }
};
int main()
{
    testArray *test = new testArray;
    test->runTest();
    //testArray test;
    //test.runTest();
    system("pause");
    return 0;
}

我想通过引用传递所有内容，因此无需复制。

__global__ void add1(int height, int width, int *&a, int *&resultArray)

通过引用传递指针有一个用途：修改（重新拔插）调用方作用域中的指针。你不这样做。因此，在这种情况下，参考文献是多余的。事实上，这是一种悲观，因为你正在引入另一个层次的间接性。请改用以下签名：

__global__ void add1(int height, int width, int* a, int* resultArray)
<小时 />

这将编译并运行，但似乎 +1 操作从未发生过。这是为什么呢？

我知道我应该有 catch 错误语句，这段代码只是一个简单的例子。

好吧，这真的很不幸，因为添加适当的错误检查可能会帮助您找到错误。将来，请考虑在询问 SO 之前添加错误检查。

内核希望其参数位于可以访问的地址空间中。这意味着它必须是通过调用任何cudaMalloc变体获得的指针。

但是你在通过什么？

myArray = new int[height*width]; // Not a cudaMalloc* variant
[...]
add1<<<dimGridWidth,dimBlock>>>(height, width, myArray, myArray);

因此，您传递给内核的指针没有任何意义，因为它不在"CUDA 地址空间"中。您的内核可能会立即出现段错误。

<小时 />

我认为您的困惑可能是由于myArray（testArray）的封闭类继承自Managed的事实引起的。这意味着new testArray将在 GPU 可访问的地址空间中分配testArray，但这并不意味着在该类成员上使用 operator new 也会在该地址空间中分配它们。它们也需要通过cudaMalloc*进行分配（例如，尽管不是必需的，但通过将分配转发到cudaMallocManaged的重载operator new）。一个简单的解决方案是不是使用 new 而是像这样分配数组：

cudaMallocManaged(&myArray, width * height* sizeof(*myArray));

将相应的 delete调用替换为 cudaFree 。

此外：

testArray test;

这不会在 GPU 可访问空间上分配test，因为它不是通过 operator new 分配的。

相关文章：