0

使用 OpenCL，我似乎无法从 Radeon 7970 将超过 7MB/秒的数据提取到 i5 Desktop 的主内存中。

#include <iostream>
#include <Windows.h>
#include <CL/cl.h>

int main(int argc, char ** argv)
{
    cl_platform_id platform;
    clGetPlatformIDs(1, &platform, NULL);
    cl_device_id device;
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
    cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);
    const char *source =
    "__kernel void copytest(__global short* dst) {\n"
    "    __local short buff[1024];\n"
    "    for (int i = 0; i < 1024; i++) {\n"
    "        for (int j = 0; j < 1024; j++)\n"
    "            buff[j] = j;\n"
    "        (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);\n"
    "    }\n"
    "}\n";
    cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
    clBuildProgram( program, 1, &device, NULL, NULL, NULL);
    cl_kernel kernel = clCreateKernel( program, "copytest", NULL);
    cl_mem buf = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 1024 * 1024 * 2, NULL, NULL);
    const size_t global_work_size = 1;
    clSetKernelArg(kernel, 0, sizeof(buf), (void*)&buf);
    LARGE_INTEGER pcFreq = {}, pcStart = {}, pcEnd = {};
    QueryPerformanceFrequency(&pcFreq);
    QueryPerformanceCounter(&pcStart);
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
    clFinish(queue);
    QueryPerformanceCounter(&pcEnd);
    std::cout << 2.0 * pcFreq.QuadPart / (pcEnd.QuadPart-pcStart.QuadPart) << "MB/sec";
}

如您所见，它只在一个工作单元上运行。我尝试用分布在多个 (64) 工作单元之间的循环替换 async_work_group_copy() ，但这没有帮助。

有什么方法可以让 Radeon 的内存速度超过 7MB/秒？我对数百 MB/秒感兴趣。英伟达会更快吗？

4

3 回答 3

1

问题是您只在 GPU 上使用一个线程，导致数千个线程处于空闲状态。在这种情况下，您可以做两件事来帮助您实现更快的速度。

首先，尝试在工作组中使用更多线程：

__kernel void copytest(__global short* dst) {
    __local short buff[1024];
    for (int i = 0; i < 1024; i++) {
        for (int j = get_local_id(0); j < 1024; j+= get_local_size(0))
            buff[j] = j;
        barrier(CLK_LOCAL_MEM_FENCE);
        (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);
    }
}

然后您可以将工作组的大小增加到 256 个左右。

const size_t local_work_size = 256;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

其次，您使用的是 GPU，因此您可能不应该只使用一个工作组。您可以使用更多工作组，如：

__kernel void copytest(__global short* dst) {
    __local short buff[1024];
    for (int i = get_group_id(0); i < 1024; i += get_num_groups(0)) {
        for (int j = get_local_id(0); j < 1024; j+= get_local_size(0))
            buff[j] = j;
        barrier(CLK_LOCAL_MEM_FENCE);
        (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);
    }
}

然后你可以增加工作组的数量：

const size_t local_work_size = 256;
const size_t global_work_size = local_work_size * 32;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

希望这将有助于加快您的应用程序。

于 2012-10-17T17:28:29.187 回答
0

确保正确分配缓冲区：阅读 NVIDIA OpenCL 编程指南（如果可以找到）以查看分配固定内存。有一个成熟的例子可以让你达到 6GB/s——同样的原则也适用于 AMD。特别是 CL_MEM_ALLOC_HOST_PTR 标志。

于 2012-10-18T21:03:19.143 回答
0

你在这里并不需要你的'j'循环。async_work_group_copy 双向工作；您可以复制到/从全局和本地空间。

//kernel will copy 2MB of short* in memory
__kernel void copytest(__global short* dst) {
  __local short buff[1024];
  for (int i = get_group_id(0); i < 1024; i += get_num_groups(0)) {
    (void)async_work_group_copy(buff, &dst[i*1024], 1024, 0);
    (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);
  }
}

opencl 1.0 规范说必须有 16kb 或更多的本地内存可供您使用（32kb+ 与 ocl v1.1 或更高版本）。许多设备实际上有 32kb。我建议轮询系统并尽可能多地使用。实际上，您仍然需要为其他目的保存一些本地内存。参见 clGetDeviceInfo (CL_DEVICE_LOCAL_MEM_SIZE)

//using 16kb local memory per work group to copy 2MB...
__kernel void copytest(__global short* dst) {
  __local short buff[16384];
  for (int i = get_group_id(0); i < 64; i += get_num_groups(0)) {
    (void)async_work_group_copy(buff, &dst[i*16384], 16384, 0);
    (void)async_work_group_copy(&dst[i*16384], buff, 16384, 0);
  }
}

如果要做的复制量足够小，您可以通过使用完成工作所需的工作组数量来摆脱“i”循环。这会导致程序集中的分支更少。

//using 32kb and 16 work groups to copy 2MB...
__kernel void copytest(__global short* dst) {
  __local short buff[32768];
  int i = get_group_id(0);
  (void)async_work_group_copy(buff, &dst[i*32768], 32768, 0);
  (void)async_work_group_copy(&dst[i*32768], buff, 32768, 0);
}
于 2012-10-19T01:55:17.410 回答