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 上的结构，因此需要创建一个指针数组，然后为每个指针创建第二个数组或指针并将指针“附加”到它的父级。 “附加”意味着它将在父数组中的适当位置填充设备指针，但必须启动内核来执行此操作。更糟糕的是，它需要再次遍历结构以创建第三维并再次附加指针。第三个维度也被填充以对齐，这是额外内存的来源。这会导致大量额外开销，并将严重影响您的性能。

此外，由于主机和设备之间的数据传输只能在连续的块上完成，运行时必须遍历结构并一次仅复制 2 个元素。再次，导致高开销。

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

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

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