基于Cuda统一内存的多维数组分配

Question

我试图通过在Power 8系统上使用CUDA UMA来分配多维数组。然而，我有问题，而尺寸越来越大。我使用的代码如下。当尺寸是24 x 24 x 24 x 5工作良好。当我将其增加到64 x 64 x 64 x 8时，即使我的设备中有内存，我也会出现"内存不足"。Afaik，我想能够通过UMA分配内存和GPU设备的物理内存一样多。所以我不希望有任何误差。目前我的主要配置是Power 8和Tesla k40，在运行时我有隔离故障。但是，我在x86 + k40机器上试用了我提供的代码片段。这招出人意料地奏效了。

顺便说一句，如果你告诉我另一种方法，除了把我所有的代码从4d数组转换成1d数组，我会很感激的。

Thanks in advance

Driver: Nvidia 361

#include <iostream>
#include <cuda_runtime.h>
void* operator new[] (size_t len) throw(std::bad_alloc) {
void *ptr;
cudaMallocManaged(&ptr, len);
return ptr;
}
template<typename T>
T**** create_4d(int a, int b, int c, int d){
T**** ary = new T***[a];
for(int i = 0; i < a; ++i)
{
ary[i] = new T**[b];
for(int j = 0; j < b; ++j){
ary[i][j] = new T*[c];
for(int k = 0; k < c; ++k){
ary[i][j][k] = new T[d];
}
}
}
return ary;
}
int main() {
double ****data;
std::cout << "allocating..." << std::endl;
data = create_4d<double>(32,65,65,5);
std::cout << "Hooreey !!!" << std::endl;
//segfault here
std::cout << "allocating..." << std::endl;
data = create_4d<double>(64,65,65,5);
std::cout << "Hooreey !!!" << std::endl;   
return 0;
}

Answer 1

在你的交叉发布中有相当多的对话，包括对你的主要问题的回答。我将用这个答案来总结一下有什么，并具体地回答这个问题:

顺便说一句，如果你告诉我另一种方法，除了把我所有的代码从4d数组转换成1d数组，我会很感激的。

1. 你的声明之一是你正在做正确的错误检查("我正确地捕获了错误。")。你不是。CUDA运行时API调用(包括cudaMallocManaged)本身做不生成c++风格的异常，所以你的throw规范对new操作符的定义是没有意义的。CUDA运行时API调用返回一个错误代码。如果您想要进行正确的错误检查，您必须收集并处理此错误代码。如果你收集了错误代码，你可以使用它来生成一个异常，如果你愿意，和一个例子，你可能会做的是包含在规范正确的CUDA错误检查问题，作为Jared Hoberock的答案之一。由于这种疏忽，当您的分配最终失败时，您忽略了这一点，然后当您试图将这些(未)分配的区域用于后续的指针存储时，您会生成一个segfault。

2. 分配失败的最接近的原因是您实际上内存不足，正如您在交叉发布中讨论的那样。您可以通过适当的错误检查轻松地确认这一点。托管分配具有粒度，因此当您请求相对较小的分配时，实际上使用的内存比您想象的要多——您请求的小分配每个都被四舍五入到分配粒度中。分配粒度的大小因系统类型而异，因此您所操作的OpenPower系统的分配粒度比您所比较的x86系统大得多，因此您没有在x86系统上耗尽内存，而是在Power系统上耗尽内存。正如你在交叉发布中所讨论的，这很容易通过对cudaMemGetInfo的战略调用来验证。

从性能角度来看，这是一种非常糟糕的多维分配方法，原因如下:

1. 您正在创建的分配是不相交的，由指针连接。因此，要通过指针解引用来访问元素，需要3或4次这样的解引用来遍历一个下标为4的指针数组。每个解引用都涉及到设备内存访问。与在1-D(平面)分配中使用模拟的4-D访问相比，这明显要慢得多。与将4-D模拟访问转换为单个线性索引相关的算法将比通过指针跟踪遍历内存快得多。

2. 由于您正在创建的分配是不相交的，托管内存子系统不能将它们合并到单个传输中，因此，在底层，在内核启动时(可能在终止时)将发生等于前3个维度乘积的许多传输。在下一次cudaDeviceSynchronize()呼叫时)。当然，这些数据必须全部传输，但是与"平坦"分配的单个传输相比，您将执行大量非常小的传输。大量小传输的相关开销可能非常大。

3. 正如我们所看到的，分配粒度会严重影响这种分配方案的内存使用效率。应该只使用一小部分系统内存的东西最终会使用所有系统内存。

4. 在这样的分配中从"行"到"行"的连续数据上工作的操作将失败，因为分配是不相交的。例如，这样一个矩阵或这样一个矩阵的一个子矩阵不能可靠地传递给CUBLAS线性代数例程，因为该矩阵的期望将具有与之相关的内存中的行存储的连续性。

理想的解决方案是创建单个平面分配，然后使用模拟的4-D索引来创建单个线性索引。这种方法将解决上述所有4个问题。然而，它可能需要大量的代码重构。

然而，我们可以提出另一种方法，它保留4下标索引，但通过创建单个底层平面分配来解决上述第2、3和4项中的问题。

下面是一个工作示例。实际上，我们将创建2个托管分配:一个底层平面分配用于数据存储，另一个底层平面分配用于指针存储(无论维度如何)。通过一些仔细的调整工作，将这两个组合成一个分配是可能的，但是这并不是实现任何提议的好处所必需的。

基本方法在SO标签上的各种其他CUDA问题中都有涉及，但其中大多数都有主机端使用(仅)考虑，因为它们没有考虑UM。然而，UM允许我们将该方法扩展到主机端和设备端使用。我们将首先创建存储数据所需大小的单个"基本"分配。然后为指针数组创建一个分配，然后遍历指针数组，将每个指针固定为指向指针数组中的正确位置，或者指向"基"数据数组中的正确位置。

下面是一个工作示例，演示了主机和设备的使用情况，并包括适当的错误检查:

$ cat t1271.cu
#include <iostream>
#include <assert.h>
template<typename T>
T**** create_4d_flat(int a, int b, int c, int d){
T *base;
cudaError_t err = cudaMallocManaged(&base, a*b*c*d*sizeof(T));
assert(err == cudaSuccess);
T ****ary;
err = cudaMallocManaged(&ary, (a+a*b+a*b*c)*sizeof(T*));
assert(err == cudaSuccess);
for (int i = 0; i < a; i++){
ary[i] =  (T ***)((ary + a) + i*b);
for (int j = 0; j < b; j++){
ary[i][j] = (T **)((ary + a + a*b) + i*b*c + j*c);
for (int k = 0; k < c; k++)
ary[i][j][k] = base + ((i*b+j)*c + k)*d;}}
return ary;
}
template<typename T>
void free_4d_flat(T**** ary){
if (ary[0][0][0]) cudaFree(ary[0][0][0]);
if (ary) cudaFree(ary);
}

template<typename T>
__global__ void fill(T**** data, int a, int b, int c, int d){
unsigned long long int val = 0;
for (int i = 0; i < a; i++)
for (int j = 0; j < b; j++)
for (int k = 0; k < c; k++)
for (int l = 0; l < d; l++)
data[i][j][k][l] = val++;
}
void report_gpu_mem()
{
size_t free, total;
cudaMemGetInfo(&free, &total);
std::cout << "Free = " << free << " Total = " << total <<std::endl;
}
int main() {
report_gpu_mem();
unsigned long long int ****data2;
std::cout << "allocating..." << std::endl;
data2 = create_4d_flat<unsigned long long int>(64, 63, 62, 5);
report_gpu_mem();
fill<<<1,1>>>(data2, 64, 63, 62, 5);
cudaError_t err = cudaDeviceSynchronize();
assert(err == cudaSuccess);
std::cout << "validating..." << std::endl;
for (int i = 0; i < 64*63*62*5; i++)
if (*(data2[0][0][0] + i) != i) {std::cout << "mismatch at "  << i << " was " << *(data2[0][0][0] + i) << std::endl; return -1;}
free_4d_flat(data2);
return 0;
}
$ nvcc -arch=sm_35 -o t1271 t1271.cu
$ cuda-memcheck ./t1271
========= CUDA-MEMCHECK
Free = 5904859136 Total = 5975900160
allocating...
Free = 5892276224 Total = 5975900160
validating...
========= ERROR SUMMARY: 0 errors
$

指出:

1. 这仍然涉及指针跟踪效率低下。如果不去掉多个下标，我不知道有什么方法可以避免这种情况。

2. 我选择在主机和设备代码中使用2种不同的索引方案。在设备代码中，我使用普通的4下标索引来演示它的效用。在主机代码中，我使用"平面"索引，以证明底层存储是连续的，并且是连续可寻址的。