2

我正在使用 Tesla m1060 进行 GPGPU 计算。它具有以下规格：

# of Tesla GPUs 1
# of Streaming Processor Cores (XXX per processor)  240
Memory Interface (512-bit per GPU)  512-bit

当我使用 OpenCL 时，我可以显示以下板子信息：

available platform OpenCL 1.1 CUDA 6.5.14
device Tesla M1060 type:CL_DEVICE_TYPE_GPU
max compute units:30 
max work item dimensions:3
max work item sizes (dim:0):512
max work item sizes (dim:1):512
max work item sizes (dim:2):64
global mem size(bytes):4294770688 local mem size:16383

如何将 GPU 卡信息与 OpenCL 内存信息相关联？

例如：

  • “内存交互”是什么意思？它是否链接了工作项？
  • 如何将 GPU 的“240 核”与工作组/项目相关联？
  • 如何将工作组映射到它（要使用的工作组数量是多少）？

谢谢

编辑：

在以下答案之后，我仍然不清楚一件事：

我使用的内核的CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE值为 32。

但是，我的设备的CL_DEVICE_MAX_COMPUTE_UNITS值为 30。

在 OpenCL 1.1 Api 中，它是这样写的（第 15 页）：

计算单元：一个 OpenCL 设备有一个或多个计算单元。工作组在单个计算单元上执行

似乎有些东西在这里不连贯，或者我没有完全理解工作组和计算单元之间的区别。

如前所述，当我将工作组的数量设置为 32 时，程序失败并出现以下错误：

Entry function uses too much shared data (0x4020 bytes, 0x4000 max).

值 16 有效。

附录

这是我的内核签名：

// enable double precision (not enabled by default)
#ifdef cl_khr_fp64
    #pragma OPENCL EXTENSION cl_khr_fp64 : enable
#else
    #error "IEEE-754 double precision not supported by OpenCL implementation."
#endif

#define BLOCK_SIZE 16 // --> this is what defines the WG size to me

__kernel __attribute__((reqd_work_group_size(BLOCK_SIZE, BLOCK_SIZE, 1)))
  void mmult(__global double * A, __global double * B, __global double * C, const unsigned int q)
{
  __local double A_sub[BLOCK_SIZE][BLOCK_SIZE];
  __local double B_sub[BLOCK_SIZE][BLOCK_SIZE];
  // stuff that does matrix multiplication with __local
}

在主机代码部分：

#define BLOCK_SIZE 16 
...
const size_t local_work_size[2] = {BLOCK_SIZE, BLOCK_SIZE};
...
status = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
4

2 回答 2

2

内存接口对 opencl 应用程序没有任何意义。它是内存控制器用于读取/写入内存的位数（现代 gpus 中的 ddr5 部分）。最大全局内存速度的公式大约是：pipelineWidth * memoryClockSpeed，但由于 opencl 是跨平台的，除非您试图找出内存性能的上限，否则您实际上不需要知道这个值。在处理内存合并时，了解 512 位接口会有些用处。wiki：合并（计算机科学）

最大工作项大小与 1) 硬件如何安排计算以及 2) 设备上的低级内存量有关 - 例如。私有内存和本地内存。

240这个数字对opencl也没有多大关系。对于此 GPU 架构，您可以确定 30 个计算单元中的每一个都由 8 个流处理器内核组成（因为 240/30 = 8）。如果您查询 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE，它很可能是该设备的 8 的倍数。请参阅：clGetKernelWorkGroupInfo

我已经回答了关于工作组规模的类似问题。看这里和这里

最终，您需要根据自己的基准测试结果调整您的应用程序和内核。我发现花时间编写许多具有不同工作组规模的测试并最终硬编码最佳规模是值得的。

于 2014-11-21T15:24:51.413 回答
1

添加另一个答案来解决您的本地内存问题。

入口函数使用过多的共享数据（0x4020 字节，最大 0x4000）

由于您正在分配 A_sub 和 B_sub，每个都有 32*32*sizeof(double)，因此您的本地内存不足。该设备应该允许您毫无问题地分配 16kb 或 0x4000 字节的本地内存。

0x4020 是 32 字节或比您的设备允许的多 4 倍。我能想到的只有两件事可能导致错误：1）您的设备或驱动程序可能存在错误，阻止您分配完整的 16kb，或 2）您正在内核中的其他位置分配内存。

您现在必须使用小于 32 的 BLOCK_SIZE 值来解决此问题。

不过也有好消息。如果您只想将 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 的倍数作为工作组大小，BLOCK_SIZE=16 已经为您完成了这项工作。（16 * 16 = 256 = 32 * 8）。为了更好地利用本地内存，请尝试 BLOCK_SIZE=24。(576=32*18)

于 2014-11-26T19:20:21.887 回答