OpenCl内核中的for循环遍历全局内存浮点数组答案

【问题标题】：For loop in OpenCl kernel rolling through global memory float arrayOpenCl内核中的for循环遍历全局内存浮点数组
【发布时间】：2017-11-17 02:06:50
【问题描述】：

我觉得我不懂基本的并行编程概念。下面的内核是一个简单/人为的示例，它重现了我遇到的问题。它尝试使用“点”中的所有值来计算一个值并将其分配给“块”中的所有项目。我想突破这些数组大小的限制。虽然我可以使“块”数组退出大（> 1 亿个浮点数），但当“点”填充超过 ~10 万个浮点数时（在 clEnqueueNDRangeKernel 之后立即调用 clFinish 之后），我得到一个“无效的命令队列”错误。谁能帮我理解为什么？

__kernel void openClTesting (__global float *blocks, __global float *points, int pointsCount)
    {
        int globalId = get_global_id(0);
        int count = 0;
        for (int i = 0; i < pointsCount; i++)
        {
            count++;
        }
        blocks[globalId] = count;
    };

一些设备信息：

CL_DEVICE_LOCAL_MEM_SIZE = 49,152  
CL_DEVICE_GLOBAL_MEM_SIZE = 2,147,483,648  
CL_DEVICE_MAX_MEM_ALLOC_SIZE = 536,870,912

主机代码：

#include "stdafx.h"
#include "CL\opencl.h"
#include <iostream>
#include <fstream>
#include <string>
#include <stddef.h>
#include <stdlib.h>
#include <stdio.h>

#define NUM_POINTS 100000
#define NUM_BLOCKS 100000000

struct openClData
{
cl_device_id deviceId = NULL;
cl_uint numDevices;
cl_uint numPlatforms;
cl_int ret;
cl_platform_id *platforms = NULL;
cl_context context;
cl_command_queue commandQueue;
cl_program program;
cl_kernel kernel;
char* kernelCode;
cl_uint kernelCodeSize;
size_t globalItemSize;
size_t localItemSize = 1;
};


char* getKernelCode();
void printErrorLog(openClData oclData);
void printRet(openClData oclData, int line);
int countFileChars(const char *fileName);

int _tmain(int argc, _TCHAR* argv[])
{
openClData oclData;
oclData.globalItemSize = NUM_POINTS;
oclData.kernelCode = getKernelCode();
std::cout << oclData.kernelCode << std::endl;
oclData.kernelCodeSize = strlen(oclData.kernelCode);

int numPoints = NUM_POINTS;
int numBlocks = NUM_BLOCKS;
cl_long localMemSize = 0, globalMemSize = 0, maxAllocMemSize = 0;
float *blocks = new float[numBlocks]{0};
float *points = new float[numPoints]{0};

//prepare platform, device, context and command queue
oclData.ret = clGetPlatformIDs(0, NULL, &oclData.numPlatforms);
printRet(oclData, __LINE__);
oclData.platforms = (cl_platform_id *)malloc(oclData.numPlatforms * sizeof(cl_platform_id));
oclData.ret = clGetPlatformIDs(oclData.numPlatforms, oclData.platforms, NULL);
printRet(oclData, __LINE__);
oclData.ret = clGetDeviceIDs(oclData.platforms[0], CL_DEVICE_TYPE_GPU, 1, &oclData.deviceId, &oclData.numDevices);
printRet(oclData, __LINE__);
oclData.context = clCreateContext(NULL, 1, &oclData.deviceId, NULL, NULL, &oclData.ret);
printRet(oclData, __LINE__);
oclData.commandQueue = clCreateCommandQueue(oclData.context, oclData.deviceId, 0, &oclData.ret);
printRet(oclData, __LINE__);
//prepare cl_mem objects
cl_mem memObjBlocks = clCreateBuffer(oclData.context, CL_MEM_READ_WRITE, sizeof(float) * numBlocks, NULL, &oclData.ret);
printRet(oclData, __LINE__);
cl_mem memObjPoints = clCreateBuffer(oclData.context, CL_MEM_READ_WRITE, sizeof(float) * numPoints, NULL, &oclData.ret);
printRet(oclData, __LINE__);
oclData.ret = clEnqueueWriteBuffer(oclData.commandQueue, memObjBlocks, CL_TRUE, 0, sizeof(float) * numBlocks, blocks, 0, NULL, NULL);
printRet(oclData, __LINE__);
oclData.ret = clEnqueueWriteBuffer(oclData.commandQueue, memObjPoints, CL_TRUE, 0, sizeof(float) * numPoints, points, 0, NULL, NULL);
printRet(oclData, __LINE__);
//prepare program
oclData.program = clCreateProgramWithSource(oclData.context, 1, (const char**)&oclData.kernelCode, (const size_t *)&oclData.kernelCodeSize, &oclData.ret);
printRet(oclData, __LINE__);
oclData.ret = clBuildProgram(oclData.program, 1, &oclData.deviceId, NULL, NULL, NULL);
printRet(oclData, __LINE__);
if (oclData.ret == CL_BUILD_PROGRAM_FAILURE) printErrorLog(oclData);
oclData.kernel = clCreateKernel(oclData.program, "openClTesting", &oclData.ret);
printRet(oclData, __LINE__);
//set arguments
oclData.ret = clSetKernelArg(oclData.kernel, 0, sizeof(cl_mem), &memObjBlocks);
printRet(oclData, __LINE__);
oclData.ret = clSetKernelArg(oclData.kernel, 1, sizeof(cl_mem), &memObjPoints);
printRet(oclData, __LINE__);
oclData.ret = clSetKernelArg(oclData.kernel, 2, sizeof(int), &numPoints);
printRet(oclData, __LINE__);
//run
oclData.ret = clEnqueueNDRangeKernel(oclData.commandQueue, oclData.kernel, 1, NULL, &oclData.globalItemSize, &oclData.localItemSize, 0, NULL, NULL);
printRet(oclData, __LINE__);
oclData.ret = clFinish(oclData.commandQueue);
printRet(oclData, __LINE__);
oclData.ret = clEnqueueReadBuffer(oclData.commandQueue, memObjBlocks, CL_TRUE, 0, sizeof(float) * numBlocks, blocks, 0, NULL, NULL);
printRet(oclData, __LINE__);
oclData.ret = clFinish(oclData.commandQueue);
printRet(oclData, __LINE__);
//print some device info
oclData.ret = clGetDeviceInfo(oclData.deviceId, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemSize, 0);
std::cout << "CL_DEVICE_LOCAL_MEM_SIZE = " << localMemSize << '\n';
oclData.ret = clGetDeviceInfo(oclData.deviceId, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_long), &globalMemSize, 0);
std::cout << "CL_DEVICE_GLOBAL_MEM_SIZE = " << globalMemSize << '\n';
oclData.ret = clGetDeviceInfo(oclData.deviceId, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_long), &maxAllocMemSize, 0);
std::cout << "CL_DEVICE_MAX_MEM_ALLOC_SIZE = " << maxAllocMemSize << '\n';

//clean up
oclData.ret = clFlush(oclData.commandQueue);
printRet(oclData, __LINE__);
oclData.ret = clFinish(oclData.commandQueue);
printRet(oclData, __LINE__);
oclData.ret = clReleaseKernel(oclData.kernel);
printRet(oclData, __LINE__);
oclData.ret = clReleaseProgram(oclData.program);
printRet(oclData, __LINE__);
oclData.ret = clReleaseMemObject(memObjBlocks);
printRet(oclData, __LINE__);
oclData.ret = clReleaseMemObject(memObjPoints);
printRet(oclData, __LINE__);
oclData.ret = clReleaseCommandQueue(oclData.commandQueue);
printRet(oclData, __LINE__);
oclData.ret = clReleaseContext(oclData.context);
printRet(oclData, __LINE__);
for (size_t i = 0; i < 10; i++)
{
    std::cout << blocks[i] << std::endl;
}
delete blocks;
delete points;
return 0;
}

char* getKernelCode()
{
char* kernelCode =
    "__kernel void openClTesting (__global float *blocks, __global float *points, int pointsCount)"
    "{"
    "   int globalId = get_global_id(0);"
    "   int count = 0;"
    "   for (int i = 0; i < pointsCount; i++)"
    "   {"
    "       count++;"
    "   }"
    "blocks[globalId] = count;"
    "}";
return kernelCode;
}

void printErrorLog(openClData oclData)
{
size_t log_size;
clGetProgramBuildInfo(oclData.program, oclData.deviceId, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
char *log = (char *)malloc(log_size);
clGetProgramBuildInfo(oclData.program, oclData.deviceId, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
std::cout << log;
free(log);
}

void printRet(openClData oclData, int line)
{
std::cout << line << ", " << oclData.ret << std::endl;
}

int countFileChars(const char *fileName)
{
std::ifstream ifs(fileName);
ifs.seekg(0, std::ios_base::end);
size_t count = ifs.tellg();
ifs.seekg(0, std::ios_base::beg);
return count;
}

【问题讨论】：

您是否查看过您的 OpenCL 实现例如为 CL_DEVICE_GLOBAL_MEM_SIZE and CL_DEVICE_MAX_MEM_ALLOC_SIZE 返回的内容，以及这些值与您遇到的明显限制之间的关系？关于“无效命令队列”错误，你能发布你的主机代码吗？（缓冲区分配、排队等）如果没有这些信息，我怀疑你会得到很多帮助。
感谢您抽出宝贵时间。我已经更新了帖子。

标签： opencl

【解决方案1】：

我注意到一些事情：

您正在启动NUM_POINTS 工作项，但将每个结果写入blocks[globalId] - 其中有NUM_BLOCKS 项。因此，当 NUM_POINTS 大于 NUM_BLOCKS 时，这是未定义的行为。它还解释了为什么改变 NUM_BLOCKS 什么都不做（在上述限制之外）：除了内存分配之外，NUM_BLOCKS 的值没有任何影响。（并且您发现的内存分配限制与您的实现的 CL_DEVICE_MAX_MEM_ALLOC_SIZE 值大致匹配。）
您可能会在此处遇到内核超时情况。在单个工作项中进行 100000 次循环迭代是相当多的。根据 OpenCL 的实现，如果内核运行时间过长，它们可能会被终止。考虑更好地利用可用的线程并行性，并在工作项之间更水平地拆分工作，而不是循环。许多短期运行的工作项通常比少数长期运行的工作项要好。

【讨论】：

【解决方案2】：

一般来说，应该避免使用localItemSize = 1;，因为它会强制每个 OpenCL 工作组由单个工作项组成，这会降低您的计算设备可以并行运行的工作组数量的并行度，这将远远少于它可以运行的工作项的数量。您可以简单地将 NULL 传递给本地项目大小，而不是让 OpenCL 实现自己找出一个合理的值：

clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, NULL, 0, NULL, NULL);

这也可能是您的错误来源，因为您正在创建 NUM_POINTS 工作组，但设备上的队列大小受内存限制 (CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE)。

【讨论】：