OpenCL如何獲取最小線程並行粒度

由於OpenCL是為各類處理器設備而打造的開發標准的計算語言。因此跟CUDA不太一樣的是，其對設備特征查詢的項更上層，而沒有提供一些更為底層的特征查詢。比如，你用OpenCL的設備查詢API只能獲取最大work group size，但無法獲取到最小線程並行粒度。

但是，由於最小線程並行粒度對於OpenCL應用領域最廣的GPU而言確實是一個比較重要的參數。如果你的work group的work item的個數是最小線程並行粒度的倍數，那么你的OpenCL kernel程序往往會達到很高的計算效率，同時也能基於這個模型來做一些Memory Bank Confliction的避免措施。因此，我這里提供了一個比較簡單的OpenCL kernel來獲取當前GPU或其它處理器的最小線程並行粒度。

我們知道，一個計算設備由若干個Compute Unit構i成，而一個Compute Unit中包含了多個Processing Element，一個Compute Unit中的所有Processing Element對於一條算術邏輯指令而言是同時進行操作的。而不同的Compute Unit之間也可以是同時進行操作。因此，GPU的並行可以划分為兩個層次——一層是Compute Unit內的所有Processing Element的並行操作；另一層是各個Compute Unit的並行操作。

上面是物理層面，如果對於OpenCL邏輯層面，我們可以認為，一個work group的最大work item個數是指一個compute unit最多能調度、分配的線程數。這個數值一般就是一個CU內所包含的PE的個數的倍數。比如，如果一個GPU有2個CU，每個CU含有8個PE，而Max work group size是512，那么說明一個CU至少可以分配供512個線程並發操作所需要的各種資源。由於一個GPU根據一條算術邏輯指令能對所有PE發射若干次作為一個“原子的”發射操作，因此，這一個對程序員而言作為“原子的”發射操作啟動了多少個線程，那么我們就可以認為是該GPU的最小並行線程數。如果一款GPU的最小線程並行數是32，那么該GPU將以32個線程作為一組原子的線程組。這意味着，如果遇到分支，那么一組32個線程組中的所有線程都將介入這個分支，對於不滿足條件的線程，則會等到這32個線程中其它線程都完成分支處理之后再一起執行下面的指令。

如果我將work group size指定為64，並且在kernel程序里加一個判斷，如果pid小於32做操作A，否則做操作B，那么pid為0~31的線程組會執行操作A，而pid為32到63的線程組不會受到阻塞，而會立馬執行操作B。此時，兩組線程將並發操作（注意，這里是並發，而不是並行。因為上面講過，GPU一次發射32個線程的話，那么對於多個32線程組將會調度發射指令）。

根據這個特性，我們就可以寫一個OpenCL kernel程序來判別當前GPU的最小並行線程粒度。

我們首先會將work group size定為最大能接受的尺寸。然后，我們將這個work group平均划分為兩組，對它們進行測試。我們在中間定義了一個local memory的變量，每個線程都能訪問它，不過我們只讓pid為0以及pid為[max_work_group_size / 2]的線程去訪問它，以不受太多干擾。如果這個標志在線程組0執行時被線程組1改變，那么我們就知道這個粒度並非是最小的，然后對前一組再平均划分為2，遞歸操作。如果在執行線程組0之后標志沒有被更改，那么說明這整個線程組是一個原子的線程組，也就是我們所要的最小並行的線程粒度。

在內核程序中，我們還傳了一個用於延遲的循環次數，使得非原子的線程組能夠被並發執行。

下面的程序的執行環境為：Windows 7 32-bit Home Edition    AMD-APU A6-3420M    Visual Studio 2013 Express Edition    AMD APP SDK

下面先貼主機端的部分代碼片斷：

/*Step 3: Create context.*/
        cl_context context = nullptr;       // OpenCL context
        cl_command_queue commandQueue = nullptr;
        cl_program program = nullptr;       // OpenCL kernel program object that'll be running on the compute device
        cl_mem outputMemObj = nullptr;      // output memory object for output
        cl_kernel kernel = nullptr;         // kernel object
        const int deviceIndex = 0;

        context = clCreateContext(NULL,1, &devices[deviceIndex],NULL,NULL,NULL);

        /*Step 4: Creating command queue associate with the context.*/
        commandQueue = clCreateCommandQueue(context, devices[deviceIndex], 0, NULL);

        /*Step 5: Create program object */
        // Read the kernel code to the buffer
        FILE *fp = fopen("cl_kernel.cl", "rb");
        if(fp == nullptr)
        {
            puts("The kernel file not found!");
            goto RELEASE_RESOURCES;
        }
        fseek(fp, 0, SEEK_END);
        size_t kernelLength = ftell(fp);
        fseek(fp, 0, SEEK_SET);
        char *kernelCodeBuffer = (char*)malloc(kernelLength + 1);
        fread(kernelCodeBuffer, 1, kernelLength, fp);
        kernelCodeBuffer[kernelLength] = '\0';
        fclose(fp);
        
        const char *aSource = kernelCodeBuffer;
        program = clCreateProgramWithSource(context, 1, &aSource, &kernelLength, NULL);

        /*Step 6: Build program. */
        status = clBuildProgram(program, 1, &devices[deviceIndex], NULL, NULL, NULL);

        /*Step 7: Initial inputs and output for the host and create memory objects for the kernel*/
        cl_int outputArg = 0;
        outputMemObj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(outputArg), NULL, NULL);

        /*Step 8: Create kernel object */
        kernel = clCreateKernel(program,"QueryMinimumGranularity", NULL);

        /*Step 9: Sets Kernel arguments.*/
        cl_int inputArg = 1000;
        status = clSetKernelArg(kernel, 0, sizeof(inputArg), &inputArg);
        status = clSetKernelArg(kernel, 1, sizeof(outputMemObj), &outputMemObj);

        /*Step 10: Running the kernel.*/
        size_t groupSize;
        clGetDeviceInfo(devices[deviceIndex], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(groupSize), &groupSize, NULL);
        size_t global_work_size[1] = { groupSize };
        size_t local_work_size[1] = { groupSize };
        status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
        clFinish(commandQueue);     // Force wait until the OpenCL kernel is completed

        /*Step 11: Read the cout put back to host memory.*/
        status = clEnqueueReadBuffer(commandQueue, outputMemObj, CL_TRUE, 0, sizeof(outputArg), &outputArg, 0, NULL, NULL);
        char chBuffer[256];
        wchar_t wsBuffer[256];
        sprintf(chBuffer, "The minimum granularity is: %d", outputArg);
        MBString2WCString(wsBuffer, chBuffer, false);
        MessageBox(hWnd, wsBuffer, L"Notice", MB_OK);

下面是kernel代碼：

__kernel void QueryMinimumGranularity(int nLoop, __global int *pOut)
{
    __local volatile int flag;

    int index = get_global_id(0);
    int totalItems = get_global_size(0);

    do
    {
        int halfIndex = totalItems / 2;
        if(index == 0)
            flag = 1;

        barrier(CLK_LOCAL_MEM_FENCE);

        if(index < halfIndex)
        {
            for(int i = 0; i < nLoop; i++)
            {
                if(flag == -1)
                    break;
            }
            if(flag != -1)
            {
                if(index == 0)
                {
                    *pOut = totalItems;
                    flag = 2;
                }
            }
        }
        else
        {
            if(index == halfIndex)
            {
                if(flag != 2)
                {
                    //while(flag != 1);
                    flag = -1;
                }
            }
        }

        barrier(CLK_LOCAL_MEM_FENCE);

        if(flag == 2)
            break;

        totalItems /= 2;
    }
    while(totalItems > 0);
}

 

 對於Windows 7小如何做基於AMD APU的OpenCL的開發，可以參考這個貼：

http://www.cnblogs.com/zenny-chen/archive/2013/06/14/3136158.html