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。