c++ - 使用结构作为缓冲支架

Question

在我当前的 OpenCL 实现中，我想节省参数时间，避免每次我想在内核中使用缓冲区时都传递它们，并为我的内核提供更短的参数列表。

因此，我创建了一个结构(工作空间)，用于保存指向设备内存中缓冲区的指针，该结构就像一个带有成员变量的对象，您希望随时间访问该对象，并且希望在整个执行过程中保持事件状态。我在 AMD GPU 甚至 CPU 上从来没有遇到过问题。但英伟达在这方面造成了很多问题。它似乎总是一个对齐问题，永远不会到达正确的缓冲区等。

这里有一些代码可以提供帮助，请参阅下面的问题:

主机上定义的结构:

 #define SRC_IMG 0       // (float4 buffer) Source image
 #define LAB_IMG 1       // (float4 buffer) LAB image

 // NOTE: The size of this array should be as much as the last define + 1.
 #define __WRKSPC_SIZE__ 2 

 // Structure defined on host.
 struct Workspace
 {
      cl_ulong getPtr[__WRKSPC_SIZE__];
 };

 struct HostWorkspace
 {
      cl::Buffer srcImg;
      cl::Buffer labImg;
 };

设备上定义的结构:

typedef struct __attribute__(( packed )) gpuWorkspace
{
    ulong getPtr[__WRKSPC_SIZE__]; 
} gpuWorkspace_t; 

请注意，在设备上，我使用 ulong，在主机上我使用 cl_ulong，如下所示 OpenCL: using struct as kernel argument .

因此，一旦创建了源图像或 LAB 图像的 cl::Buffer，我会将它们保存到 HostWorkspace 对象中，因此在该对象被释放之前，对 cl::Buffer 的引用将保留，因此整个项目都存在缓冲区在主机上，事实上在设备上。

现在，我需要为这些设备提供数据，因此我有一个简单的内核，它初始化我的设备工作区，如下所示:

__kernel void Workspace_Init(__global gpuWorkspace_t* wrkspc,
                             __global float4* src,
                             __global float4* LAB)
{
    // Get the ulong pointer on the first element of each buffer.
    wrkspc->getPtr[SRC_IMG] = &src[0];
    wrkspc->getPtr[LAB_IMG] = &LAB[0];
}

其中wrkspc是用struct Workspace分配的缓冲区，而src + LAB只是作为一维数组图像分配的缓冲区。

之后，在我的任何内核中，如果我想使用 src 或 LAB，我会执行以下操作:

__kernel void ComputeLABFromSrc(__global gpuWorkspace_t* wrkSpc)
{
    // =============================================================
    // Get pointer from work space.
    // =============================================================

    // Cast back the pointer of first element as a normal buffer you
    // want to use along the execution of the kernel.
    __global float4* srcData = ( __global float4* )( wrkSpc->getPtr[SRC_IMG] );
    __global float4* labData = ( __global float4* )( wrkSpc->getPtr[LAB_IMG] );

    // Code kernel as usual.
}

当我开始使用这个时，我有大约 4-5 个图像，进展顺利，具有不同的结构，如下所示:

struct Workspace
{
    cl_ulong imgPtr;
    cl_ulong labPtr;
};

其中每个图像都有自己的指针。

在某个时刻，我获得了更 multimap 像，但遇到了一些问题。所以我在网上搜索，发现一些建议，结构的 sizeof() 在设备/主机之间可能不同，所以我将其更改为同一时间的单个数组，并且在 16 个元素之前都可以正常工作。

因此我进行了更多搜索，找到了关于属性((packed))的推荐，我将其放在设备结构上(见上文)。但现在，我达到了 26 个元素，当我检查设备或主机上的结构大小时，大小为 208 (elements * sizeof(cl_ulong) == 26 * 8)。但我仍然遇到与之前的模型类似的问题，我的指针在上一个图像中间的其他地方读取，等等。

所以我想知道是否有人尝试过类似的模型(可能采用不同的方法)或有任何技巧来建立一个“可靠”的模型。

请注意，所有内核都经过良好编码，在 AMD 或 CPU 上使用相同代码执行时，我得到了很好的结果。唯一的问题在于 Nvidia。

Answer 1

不要尝试跨内核边界存储 GPU 端指针值。不保证它们保持不变。始终使用索引。如果内核使用特定的缓冲区，您需要将其作为参数传递给该内核。

引用文献:

1. OpenCL 1.2 规范(据我所知，nvidia 没有实现更新的标准)没有定义指针到整数转换的行为，反之亦然。
2. 第 6.9p 节指出:“声明为结构体或 union 体的内核函数的参数不允许 OpenCL 对象作为结构体或 union 体的元素传递。”这正是您正在尝试执行以下操作:将缓冲区结构传递给内核。
3. 第 6.9a 节规定:“程序中内核函数的参数不能声明为指向指针的指针。” - 这本质上是您试图通过强制转换来破坏的内容你的指针指向一个整数并返回。 (第 1 点)您无法通过绕过类型系统来“欺骗”OpenCL 使其得到明确定义。

正如我在下面的评论线程中建议的那样，您将需要使用索引来保存缓冲区对象内的位置。如果您想跨不同的内存区域存储位置，则需要将多个缓冲区统一为一个，并将一个索引保存到这个巨大的缓冲区中，或者保存一个数字值来标识您所引用的缓冲区。