OpenCl 程序未使用 AMD 280X GPU 中的所有可用内核

Question

伙计们。

我正在 OpenCL 1.2 上开发一个序列比较应用程序，并在 AMD R9 280X GPU 中进行测试。以下是显卡信息：

DRIVER_VERSION: 1445.5 (VM)
Type: GPU
EXECUTION_CAPABILITIES: Kernel
GLOBAL_MEM_CACHE_TYPE: Read-Write (2)
CL_DEVICE_LOCAL_MEM_TYPE: Local (1)
SINGLE_FP_CONFIG: 0xbe
QUEUE_PROPERTIES: 0x2
VENDOR_ID: 4098
MAX_COMPUTE_UNITS: 32
MAX_WORK_ITEM_DIMENSIONS: 3
MAX_WORK_GROUP_SIZE: 256
PREFERRED_VECTOR_WIDTH_CHAR: 4
PREFERRED_VECTOR_WIDTH_SHORT: 2
PREFERRED_VECTOR_WIDTH_INT: 1
PREFERRED_VECTOR_WIDTH_LONG: 1
PREFERRED_VECTOR_WIDTH_FLOAT: 1
PREFERRED_VECTOR_WIDTH_DOUBLE: 1
MAX_CLOCK_FREQUENCY: 1020
ADDRESS_BITS: 32
MAX_MEM_ALLOC_SIZE: 1073741824
IMAGE_SUPPORT: 1
MAX_READ_IMAGE_ARGS: 128
MAX_WRITE_IMAGE_ARGS: 8
IMAGE2D_MAX_WIDTH: 16384
IMAGE2D_MAX_HEIGHT: 16384
IMAGE3D_MAX_WIDTH: 2048
IMAGE3D_MAX_HEIGHT: 2048
IMAGE3D_MAX_DEPTH: 2048
MAX_SAMPLERS: 16
MAX_PARAMETER_SIZE: 1024
MEM_BASE_ADDR_ALIGN: 2048
MIN_DATA_TYPE_ALIGN_SIZE: 128
GLOBAL_MEM_CACHELINE_SIZE: 64
GLOBAL_MEM_CACHE_SIZE: 16384
GLOBAL_MEM_SIZE: 2893021184
MAX_CONSTANT_BUFFER_SIZE: 65536
MAX_CONSTANT_ARGS: 8
LOCAL_MEM_SIZE: 32768
ERROR_CORRECTION_SUPPORT: 0
PROFILING_TIMER_RESOLUTION: 1
ENDIAN_LITTLE: 1
AVAILABLE: 1
COMPILER_AVAILABLE: 1
MAX_WORK_GROUP_SIZES: 256 256 256

程序是正确的并且产生了正确的结果（它也可以在其他 CPU 和 GPU 处理器上运行），但是性能很差。在我看来，OpenCL 并没有使用所有可用的内核。相同的代码在 Nvidia GTX 680 卡上运行速度提高 50 倍。

代码有点复杂，所以我只发布主机代码，一旦 OpenCL 代码正确执行。

    err = 0;
    err  = clSetKernelArg(kernel2, 0, sizeof(i0), &i0);
    err |= clSetKernelArg(kernel2, 1, sizeof(i1), &i1);
    err |= clSetKernelArg(kernel2, 2, sizeof(step), &step);
    err |= clSetKernelArg(kernel2, 3, sizeof(cutBlock), &cutBlock);
    err |= clSetKernelArg(kernel2, 4, sizeof(cl_mem), (void*) &op->d_blockResult);
    err |= clSetKernelArg(kernel2, 5, sizeof(cl_mem), (void*) &op->d_busH);
    err |= clSetKernelArg(kernel2, 6, sizeof(cl_mem), (void*) &op->d_extraH);
    err |= clSetKernelArg(kernel2, 7, sizeof(cl_mem), (void*) &op->d_busV_h);
    err |= clSetKernelArg(kernel2, 8, sizeof(cl_mem), (void*) &op->d_busV_e);
    err |= clSetKernelArg(kernel2, 9, sizeof(cl_mem), (void*) &op->d_busV_o);
    err |= clSetKernelArg(kernel2, 10, sizeof(cl_mem), (void*) &op->d_split_m);
    err |= clSetKernelArg(kernel2, 11, sizeof(cl_mem), (void*) &op->t_seq0);
    err |= clSetKernelArg(kernel2, 12, sizeof(cl_mem), (void*) &op->t_seq1);

    if (err != CL_SUCCESS)
        exit(0);

    global = blocks * threads;
    local = threads;

    err = clGetKernelWorkGroupInfo(kernel2, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_work_groups), &max_work_groups, NULL);
    if (err != CL_SUCCESS)
        exit(0);

    err = clEnqueueNDRangeKernel(commands, kernel2, 1, NULL, &global, &local, 0, NULL, &k_event);
    if (err != CL_SUCCESS)
        exit(0);

    err = clWaitForEvents(1,&k_event);
    err = clReleaseEvent(k_event);

在典型的执行中，global=4096 和 local=64，所以工作量大于 GPU 核心数。

OpenCL 有什么理由不使用所有可用的内核？会不会是驱动错误？

提前致谢。

马可

Answer 1

 global=4096 and local=64, so the work size is greater than the number of GPU cores.

这还不够。您的 gpu 有 2048 个内核，但内核执行开销使执行本身相形见绌，因为​​每个内核只工作两次。根据您在 GPU 上工作的算法，您至少需要 8192,16384 甚至 1M 全局大小。当内存单元也很忙时，您需要让这些核心保持忙碌。

Nvidia gpu 可以设置 local=1024 并在 4 个 smx 块内完成整个工作。 （增加线程间通信能力）。当工作量足够大时，您的卡可能会被压倒。

您是否将缓冲区上传-togpu 和缓冲区-从-gpu 下载时间排除在性能计算之外？

Any reason to OpenCL not to use all available cores? 

由于 opencl 没有告诉我们每个内核的使用情况，如果你的 gpu 发热与 gtx680 一样多，你只能说“没关系”，因为相信“热源是内核的半导体”。

The same code runs 50X faster in a Nvidia GTX 680 card.

相同的代码，但 CUDA？开放式？与互操作？你自己都测试了吗？使用 sse/avx 对抗单核 CPU 或所有核或所有核？相同的驱动版本？同一个操作系统？相同的环境？

你的算法是什么？如果那是矩阵乘法，则需要更多的线程间通信，请使用 local=256(nvidia 有 1024) 您是否使用异步副本？您的代码有任何虚假的递归性吗？