OpenCL如何判定一个work-group的最大Local Memory大小

最近有不少朋友提及到如何能在运行时获悉一个GPU的最大local memory的尺寸。由于OpenCL对各类处理器开放，因此不同处理器所拥有的local memory大小也各不相同。即便是GPU，甚至同一家公司出的GPU，不同的架构，其Local Memory的尺寸也各不相同。一般来说，现在随着制程工艺的不断发展，Local Memory也逐步变大。

这里简单地通过运行时动态试探法来获悉你当前所使用的GPU，一个Compute Unit上的Local Memory大小。至少，我们可以了解到一个work group最多能容纳多少大小的local memory。

其实原理很简单。如果你针对一个特定的GPU写的kernel代码，里面给一个work group分配的local memory超过了其最大尺寸限制，则在构建kernel程序的时候编译器就会报错。我们通过这一特点来试探当前GPU的每个work-group最大的local memory尺寸是多少。

代码很简单，直接在主机端写OpenCL内核代码字符串。每次比前一次多申请4KB的大小，然后看看构建结果。

    /** Query local memory size */
    char kernelSrcBuffer[256];
    int localMemorySize = 0;
    
    for(int i = 1; i < 1024; i++)
    {
        size_t kernelLength = sprintf(kernelSrcBuffer, "__kernel void QueryLDSSize(void){ __local int lds[%d * 1024]; }", i);
        const char *pKernelSrc = kernelSrcBuffer;
        program = clCreateProgramWithSource(context, 1, &pKernelSrc, &kernelLength, NULL);
        if(clBuildProgram(program, 1, &device, NULL, NULL, NULL) == CL_BUILD_PROGRAM_FAILURE)
        {
            clReleaseProgram(program);
            localMemorySize = (i - 1) * 4;  // 注意，这里单位是KB
            break;
        }
        clReleaseProgram(program);
    }

这里省略了cl_dvice以及cl_context变量的声明和初始化，这些由开发者自己实现，呵呵～

在Intel HD Grapchis 5000上，Local Memory的尺寸为64KB。