由于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] = ' '; 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