OpenCL和Tesla M1060

Question

我使用特斯拉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.1Api中，它是这样写的(第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);

Answer 1

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

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

对于opencl来说，240这个数字也不重要。您可以确定30个计算单元中的每个单元都由8个流处理器核心组成，用于此gpu体系结构(因为240/30 = 8)。如果您查询CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE，该设备的密钥将是8的倍数。请参阅：clGetKernelWorkGroupInfo

我已经回答了一个关于工作组规模的类似问题。请看这里，还有这里

最终，您需要根据自己的工作台标记结果对应用程序和内核进行优化。我发现有必要花时间编写许多具有不同工作组大小的测试，并最终对最佳大小进行硬编码。

Answer 2

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

条目函数使用过多的共享数据(0x4020字节，0x4000最大)

由于您正在分配A_sub和B_sub，每个都有32*32* Since of (Double)，因此本地内存不足。设备应该允许您分配16 an，即0x4000字节的本地内存，而不会出现问题。

0x4020是32字节，比设备允许的多出4倍。我认为只有两件事可能会导致错误: 1)设备或驱动程序可能有错误，阻止您分配完整的16 of，或者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)