OpenCL Nvidia GPU 上的小型恒定内存大小答案

【问题标题】：OpenCL Small constant memory size on Nvidia GPUOpenCL Nvidia GPU 上的小型恒定内存大小
【发布时间】：2020-07-24 20:13:24
【问题描述】：

我有以下用于矩阵乘法的代码，为简单起见而缩写。我计划使用block_size*block_size 的本地内存来保存一个子矩阵块。当我在 NVIDIA GPU 上运行它时，我在clEnqueueNDRangeKernel 中不断收到错误代码-52。经过一番研究，我发现 NVIDIA gpu 上的恒定内存大小非常小。

主持人：

    cl::Buffer a_buf{ context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, a.bytes(), a.data };
    cl::Buffer b_buf{ context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, a.bytes(), bT.data };
    cl::Buffer result_buf{ context, CL_MEM_READ_WRITE , result.bytes(), nullptr }; //for memory mapping
    kernel.setArg(0, a_buf);
    kernel.setArg(1, b_buf);
    kernel.setArg(2, local_size*local_size* sizeof(float), nullptr);
    kernel.setArg(3, local_size*local_size* sizeof(float), nullptr);
    kernel.setArg(4, result_buf);
    queue.enqueueNDRangeKernel(kernel, { 0,0 }, { a.rows, a.rows }, {local_size, local_size});
                                        //  ^ offset   ^global work size  ^local work size

内核：

__kernel void matrixMul(__constant float* a,
    __constant float* b,    //storing the original matrix data
    __local float* a_local, 
    __local float* b_local, //storing a sub-matrix block for the work-group
    __global float* result)
 {...}

使用CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE，我的 RX580 会返回几乎所有可用的 VRAM，但我的 GTX1650 只返回 64KB。当使用__constant 而不是__global 时，我的RX580 确实获得了显着的性能提升。 是不是我做错了什么，或者我需要准备不同的内核才能在 AMD 和 NVIDIA gpus 上运行？

编辑：我在 github here 上发现了一个相关问题所以我更改了__constant float* a -> __global const float* restrict a，它可以工作。

【问题讨论】：

标签： opencl

【解决方案1】：

NVIDIA GPU 上的恒定内存大小确实非常小，只有 64KB（我检查了 Titan Xp、GTX 960M、RTX 2080 Ti、Tesla K20c、Tesla K40m）。在 AMD Radeon VII 上，恒定内存大小要大得多，为 14GB。在 Intel CPU（i7-8700K、Xeon E5-2680 v2）上，恒定内存大小为 128KB。这是一个驱动程序限制，解决方法是使用global const float* restrict（以及restrict 用于float* 类型的所有其他内核参数）而不是constant float*，正如您已经发现的那样。

如果 AMD GPU 的性能差异很大，您可以为 AMD 和 NVIDIA GPU 使用不同的内核声明。您可以在编译 OpenCL 代码之前通过#ifdef 或在运行时通过string 连接在它们之间切换；这样您就不必在代码中包含两次整个内核，而只需在内核参数声明的那一行。

【讨论】：