Opencl：从 __constant 复制到 __global 内存

Posted 2023-04-15

【中文标题】Opencl：从 __constant 复制到 __global 内存

【英文标题】：Opencl: Copying from __constant to __global memory

【发布时间】：2013-07-27 14:28:39

【问题描述】：

我有两个内核（它们都只运行一次，因此示例中 globalWorkSize 为 1）：

第一个内核 (kernel_Calc()) 计算一些值并将它们存储在 __global 内存中。在这个例子中，它计算（设置转换 3D 空间中的一个点的转换矩阵）一个转换矩阵并转换原点。

inline 
float4 mul( const float4 M[ 4 ], const float4 v)

   float4 r;
   r.x = dot( v, M[ 0 ] );
   r.y = dot( v, M[ 1 ] );
   r.z = dot( v, M[ 2 ] );
   r.w = dot( v, M[ 3 ] );
   return r;


__kernel
void kernel_Calc( __global float4* g_TransformationMatrices, __global float3* g_Point3D )

    __private float4 transformationMatrix[ 4 ];

    transformationMatrix [ 0 ] = (float4) ( 1.0f, 0.0f, 0.0f, 0.0f );
    transformationMatrix [ 1 ] = (float4) ( 0.0f, 1.0f, 0.0f, 10.0f );
    transformationMatrix [ 2 ] = (float4) ( 0.0f, 0.0f, 1.0f, 0.0f );
    transformationMatrix [ 3 ] = (float4) ( 0.0f, 0.0f, 0.0f, 1.0f );

    g_TransformationMatrices[ 0 ] = transformationMatrix[ 0 ];
    g_TransformationMatrices[ 1 ] = transformationMatrix[ 1 ];
    g_TransformationMatrices[ 2 ] = transformationMatrix[ 2 ];
    g_TransformationMatrices[ 3 ] = transformationMatrix[ 3 ];


    float4 point4D = (float4) ( 0.0f, 0.0f, 0.0f, 1.0f );
    float4 point4DTransformed = mul( transformationMatrix, point4D);

    g_Point3D[ 0 ] = (float3) ( point4DTransformed.x / point4DTransformed.w ,
                                point4DTransformed.y / point4DTransformed.w ,
                                point4DTransformed.z / point4DTransformed.w );

在主机端，我使用clEnqueueCopyBuffer() 函数将计算出的__global 缓冲区复制到__constant 缓冲区（CL_MEM_READ_ONLY 缓冲区）中。 （我这样做是因为我希望从__constant 内存读取比从__global 内存读取更快。缓冲区复制可以在设备端使用此功能完成，而无需将__global 复制回主机和然后将其复制到__constant。）

第二个内核 (kernel_Test()) 尝试将计算值加载到可以在主机端读取的 __global 变量 (__global float4* test) 中。 sizeStruct 是一个用户定义的结构，它只包含一个整数（这是矩阵和转换点的数量）。第二个和第三个参数是__constant内存中被clEnqueueCopyBuffer()函数填满的缓冲区。

struct sizeStruct

    int m_Size;
;

__kernel 
void kernel_Test( __constant struct sizeStruct* c_SS,
                  __constant float4* c_TransformationMatrices,
                  __constant float3* c_Points3D,
                  __global float4 *test )
                   
    test[ 0 ] = c_TransformationMatrices[ 0 ];
    test[ 1 ] = c_TransformationMatrices[ 1 ];
    test[ 2 ] = c_TransformationMatrices[ 2 ];
    test[ 3 ] = c_TransformationMatrices[ 3 ];

问题是当我运行程序时，测试变量包含以下内容：

1.000000, 0.000000, 0.000000, 0.000000
0.000000, 0.000000, 0.000000, 0.000000
0.000000, 0.000000, 0.000000, 0.000000
0.000000, 0.000000, 0.000000, 0.000000

但它应该包含：

1.000000, 0.000000, 0.000000, 0.000000
0.000000, 1.000000, 0.000000, 10.000000
0.000000, 0.000000, 1.000000, 0.000000
0.000000, 0.000000, 0.000000, 1.000000

我检查了__constant 变量（通过将它们复制到主机内存）它们包含正确的数据。该代码是我的程序的简化版本。这就是它可能包含不必要的操作和参数的原因。该示例已按照我的描述进行了测试和工作。

当我将__constant float3* c_Points3D 内核参数更改为__global float3* c_Points3D 内核参数时更有趣（但仍然使用由clEnqueueCopyBuffer() 函数填充的只读缓冲区）它工作正常。当我删除 __constant struct sizeStruct* c_SS 参数时，它也可以工作。 因此，kernel_Test 参数的地址空间似乎有问题，但问题出现在__constant -> __global 复制。

我在 nvidia geforce gtx 690 上运行程序，但我可以将设备（和平台）更改为 intel i7-3930k（使用 intel sdk）。使用 intel-i7 cpu 一切正常，无需更改内核代码。

Q1：为什么会出现这种奇怪的行为？有人知道我做错了什么吗？

Q2：使用cl_mem_read_only创建缓冲区并使用__global地址空间限定符是否合法？

【参考方案1】：

Q1：为什么会出现这种奇怪的行为？有人知道我做错了什么吗？

没有什么明显的想法，你能发布一个在 Windows 上编译的完整的最小工作示例吗？我想在 AMD Radeon GPU 上进行测试。

Q2：使用 cl_mem_read_only 创建缓冲区并使用 __global 地址空间限定符是否合法？

是的，这是合法的，尽管您需要添加 const 来指定缓冲区是只读的，请参阅 OpenCL 1.2 规范的第 6.5.1 节。

【讨论】：

感谢您的回答。我在这里上传了测试项目（VS2010）：dl.dropboxusercontent.com/u/23471668/OpenCLConstantTest.zip

您的代码在我的 AMD Radeon HD 6970 上正常运行。稍后我将在一些 Nvidia 显卡上进行测试。

NVIDIA 显卡在 OpenCL 中对向量的支持很糟糕，float4 可能是问题的根本原因。 AMD 的 FFT 实现也是如此，它们使用 float2，代码在 Intel CPU 和 AMD GPU 上运行良好，但在 NVIDIA 卡上却不行。