openACC：运行将 3D 阵列复制到设备时设备内存不足

Question

我是 openACC 的新手，我想弄清楚当我使用 copyin 数据指令（特别是在 C 中）时具体发生了什么。

我有一个 3D 阵列，正在尝试将其复制到设备上以进行一些计算。我遇到的问题是，当我进行复制时，设备内存不足，尽管据我估计数组应该只有 ~40 MB（20000 x 128 x 2 浮点数组）。我使用的是 GTX 950，当运行有超过 1 GB 的可用内存时（使用 nvidia-smi 检查）。

这是我用来测试的代码。我编译它 pgcc -acc -Minfo -o copytest copytest.c

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


float ***create_test_array( int nsamples, int nchan, int npol )
{
    int s, ch; // Loop variables
    float ***array;

    array = (float ***)malloc( nsamples * sizeof(float **) );

    for (s = 0; s < nsamples; s++)
    {
        array[s] = (float **)malloc( nchan * sizeof(float *) );

        for (ch = 0; ch < nchan; ch++)
            array[s][ch] = (float *)malloc( npol * sizeof(float) );
    }

    return array;
}



void test_copy( int nsamples, int nchan, int npol, float ***arr)
{
#pragma acc data pcopyin(arr[0:nsamples][0:nchan][0:npol])
#pragma acc kernels
    for (int pol = 0; pol < npol; pol++)
    {
        for (int ch = 0; ch < nchan; ch++)
        {
            for (int s = 0; s < nsamples; s++)
            {
                arr[s][ch][pol] = 0.0;
            }
        }

    }
}


void main()
{
    int nsamples = 10000;
    int nchan = 128;
    int npol = 2;
    float ***test_array = create_test_array( 2*nsamples, nchan, npol );

    test_copy( 2*nsamples, nchan, npol, test_array );
}

非常感谢任何见解。

Answer 1

这实际上是一个非常糟糕的 GPU 数据布局。问题是编译器必须匹配 GPU 上的结构，因此需要创建一个指针数组，然后为每个指针创建第二个数组或指针，并 "attach" 指向其父级的指针。 "attach" 表示它将在父数组中的适当位置填充设备指针，但必须启动内核才能执行此操作。更糟糕的是，它需要再次遍历结构以创建三维并再次附加指针。第三维也被填充以对齐，这是额外内存的来源。这会导致大量额外开销，并会严重影响您的性能。

另外，由于主机和设备之间的数据传输只能在连续的块上进行，运行时必须遍历结构并且一次只能复制 2 个元素。同样，造成高开销。

如果您可以交换 "nsamples" 和 "npol" 尺寸，您仍然会有一些开销，但 128+2x128 附件比 20000+20000x128 少很多。

或者，您可以使用 CUDA 统一内存 (-ta=tesla:managed) 让 CUDA 运行时为您进行数据移动。编译器不再需要创建 GPU 数据结构或附加指针。

第三种选择是将数组线性化（即使其成为一维数组），然后在循环中计算索引。

openACC：运行将 3D 阵列复制到设备时设备内存不足

openACC: Running out of device memory when copying 3D array to device

c

openacc

openACC：运行 将 3D 阵列复制到设备时设备内存不足

openACC: Running out of device memory when copying 3D array to device

c

openacc

openACC：运行将 3D 阵列复制到设备时设备内存不足