介紹

由于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之后標志沒有被更改，那么說明這整個線程組是一個原子的線程組，也就是我們所要的最小并行的線程粒度。

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

?OpenCL代碼片斷

執行環境為：Windows 7 64bit? AMD-APU A6-3420M??? Visual Studio 2008? ATI Strom APP SDK.

主機端的部分代碼片斷：

cl_context context = nullptr; // OpenCL contextcl_command_queue commandQueue = nullptr;cl_program program = nullptr; // OpenCL kernel program object that'll be running on the compute devicecl_mem outputMemObj = nullptr; // output memory object for outputcl_kernel kernel = nullptr; // kernel objectconst int deviceIndex = 0;context = clCreateContext(NULL,1, &devices[deviceIndex],NULL,NULL,NULL);commandQueue = clCreateCommandQueue(context, devices[deviceIndex], 0, NULL);// Read the kernel code to the bufferFILE *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);status = clBuildProgram(program, 1, &devices[deviceIndex], NULL, NULL, NULL);cl_int outputArg = 0;outputMemObj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(outputArg), NULL, NULL);kernel = clCreateKernel(program,"QueryMinimumGranularity", NULL);cl_int inputArg = 1000;status = clSetKernelArg(kernel, 0, sizeof(inputArg), &inputArg);status = clSetKernelArg(kernel, 1, sizeof(outputMemObj), &outputMemObj);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 completedstatus = 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); }

關于Image Engineering & Computer Vision的更多討論與交流，敬請關注本博客和新浪微博songzi_tea.

總結

以上是生活随笔為你收集整理的OpenCL最小线程并行粒度的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。