首先注明：我用的AMD的opencl，它有很多sample代碼，結合代碼來解釋這些API
Opencl 常用的API 匯總總結：
信息查詢函數
1.
cl_int clGetDeviceInfo(cl_device_id device,
cl_device_info param_name,
size_t param_value_size,
void * param_value,
size_t *param_value_size_ret )
參數說明
此函數用來查詢OpenCL設備信息。首先介紹其參數：
第一個參數device是clGetDeviceIDs的返回值。
第二個參數param_name是一個枚舉常量，標識要查詢的設備信息，具體有哪些信息稍后詳述。
第三個參數param_value_size聲明下一個參數param_value所指向的存儲空間的字節大小。這個大小要>=查詢參數大小。
第四個參數param_value指向要查詢參數返回到的存儲空間的地址。NULL時表示忽略。
第五個參數返回查詢到的參數的實際的大小。設為NULL則忽略。
下面我們來具體介紹一下可以查詢的幾個常用參數。
CL_DEVICE_TYPE:OpenCL設備類型。目前支持CL_DEVICE_TYPE_CPU,CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_DEFAULT或以上聯合。
CL_DEVICE_VENDOR_ID:一個唯一的供應商識別碼。一個唯一設備識別碼的例子，PCIe ID。
CL_DEVICE_MAX_COMPUTE_UNITS:OpenCL設備上並行計算單元數目。一個work-group只在一個compute unit上執行。該參數最小為1.
CL_DEVICE_MAX_WORK_ITEM_DEMENSIONS:數據並行執行模塊用來聲明global和localwork-item IDS的最大維度。該參數最小為3。參考clEnqueueNDRangeKernel，該函數第三個參數work_dim是聲明global work-item和work-group中的work-items（specify the global work-items and work items in the work-group.）中使用的維度的數目。work-item應該比0大，且<=CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS.
CL_DEVICE_MAX_WORK_ITEM_SIZES:數組類型，指能被在work-group的每一個維度聲明的work-item的最大數目。最小值（1,1,1）。
CL_DEVICE_MAX_WORK_GROUP_SIZE:在一個computeunit中執行一個kernel的work-group中work-item的最大數目。最小值為1。Maximum number of work-items in a work-group executing a kernel ona single compute unit, using the data parallelexecution model. (Refer to clEnqueueNDRangeKernel ). The minimum value is1.
舉例：
//Get max compute units
status = clGetDeviceInfo(
deviceId,
CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(cl_uint),
&maxComputeUnits,
NULL);
CHECK_OPENCL_ERROR(status, "clGetDeviceIDs(CL_DEVICE_MAX_COMPUTE_UNITS) failed");
2.
cl_int clGetKernelWorkGroupInfo(cl_kernelkernel ,
cl_device_iddevice ,
cl_kernel_work_group_info param_name,
size_t param_value_size,
void* param_value,
size_t*param_value_size_ret )
參數說明：
此函數返回指定到某一device上的kernel對象信息。同樣，先來看參數：
第一個參數要查詢的kernel對象。
第二個參數指定與kernel綁定的設備列表中的某一個設備。這個列表就是與kernel綁定的context對應的kernel列表。如果列表中只有一個device，此處device參數可以為NULL。
第三個參數指定要查詢參數名稱，這也是個枚舉值。
第三個參數是要指定的要查詢的參數返回的字節數，要>=返回值。
第四個參數指返回值指向內存空間的地址，若設為NULL則忽略。
第五個參數返回實際查詢到的參數的大小。
下面來說幾個重要的可查詢的參數：
CL_KERNEL_WORK_GROUP_SIZE:查詢在某一指定設備上執行一個kernel可以使用的最大work-group的size。OpenCL實現會使用資源，這就要求kernel確定work-group大小。
This provides a mechanism for the applicationto query the maximum work-group size that can be used to execute a kernel on aspecificdevice given by device. The OpenCL implementationuses the resource requirements of the kernel (register usage etc.) to determinewhat this work -group size should be.
舉例：
status = clGetKernelWorkGroupInfo(kernel,
device,
CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(cl_ulong),
&localMemoryUsed,
NULL);
if(checkVal(status, CL_SUCCESS, "clGetKernelWorkGroupInfo failed(CL_KERNEL_LOCAL_MEM_SIZE)"))
return SDK_FAILURE;
兩個例子均來自SDKCommon.cpp.
work-group/work-item/size等關系說明
為執行一個數據並行kernel，除work-items的數目外也要指定work-groups的數目。這也就是為什么兩個參數都必須傳遞給clEnqueueNDRangeKernel。例如：
size_t global_item_size = 4;//總的線程數
size_t local_item_size = 1;//每一個group的線程數
/* Execute OpenCL kernel as data parallel*/
ret = clEnqueueNDRangeKernel(command_queue,kernel, 1, NULL,
&global_item_size,&local_item_size, 0, NULL, NULL);
這個就表示上面這個數據並行計算的kernel中每一個work-group由1個work-item組成，而共有4個work-items要被處理，即總的work-items要被分成4個work-group。
另外work-item對應硬件上的一個PE（processing element）,而一個work-group對應硬件上的一個CU（computing unit）。這種對應可以理解為，一個work-item不能被拆分到多個PE上處理；同樣，一個work-group也不能拆分到多個CU上同時處理（忘了哪里看到的信息）。當映射到OpenCL硬件模型上時，每一個work-item運行在一個被稱為處理基元（processing element）的抽象硬件單元上，其中每個處理基元可以處理多個work-item(注：摘自《OpenCL異構計算》P87)。（如此而言，是不是說對於二維的globalx必須是localx的整數倍，globaly必須是localy的整數倍？那么如果我數據很大，work-item所能數量很多，如果一個group中中work-item的數量不超過CU中PE的個數，那么group的數量就可能很多；如果我想讓group數量小點，那work-item的數目就會很多，還能不能處理了呢？以當前這個示例是能的，但是對於多的work-item,這涉及到如何確定work-item數目的問題。
結合Cuda的概念進行解釋：因為實際上，一個 SM 可以允許的 block 數量，還要另外考慮到他所用到 SM 的資源：shared memory、registers 等。在 G80 中，每個 SM 有 16KB 的 shared memory 和 8192 個 register。而在同一個 SM 里的 block 和 thread，則要共享這些資源;如果資源不夠多個 block 使用的話，那 CUDA 就會減少 Block 的量，來讓資源夠用。在這種情形下，也會因此讓 SM 的 thread 數量變少，而不到最多的 768 個。
比如說如果一個 thread 要用到 16 個 register 的話(在 kernel 中宣告的變量)，那一個 SM 的 8192 個 register 實際上只能讓 512 個 thread 來使用;而如果一個 thread 要用 32 個 register，那一個 SM 就只能有 256 個 thread 了～而 shared memory 由於是 thread block 共享的，因此變成是要看一個 block 要用多少的 shread memory、一個 SM 的 16KB 能分給多少個 block 了。
所以雖然說當一個 SM 里的 thread 越多時，越能隱藏 latency，但是也會讓每個 thread 能使用的資源更少。因此，這點也就是在優化時要做取舍的了
繼續向下解釋work-group,work-item,size的關系：
每一個work-group中work-item的數目是不能改變的，始終如一。如果work-item的數目不能在work-groups中均分,clEnqueueNDRangeKernel失敗，返回錯誤碼CL_INVALID_WORK_GROUP_SIZE。此處要注意，自己在嘗試檢測GPU處理能力的時候給出的work-item和work-group的數目不能整除時不一定是數量超限，有可能只是不能整除。
global work-item ID、localwork-item ID，和work-group ID之間的關系如下圖所示。
圖1 work-group ID和work-item ID
表1 獲取ID的函數
| 函數 |
返回值 |
| get_group_id |
Work-group ID |
| get_global_id |
Global work-item ID |
| get_local_id |
Local work-item ID |
因為要處理2D圖像或3D空間，work-items和work-groups可以被指定為2或3維。圖2給出一個work-group和work-item被定義為2D的例子。
圖2 work-group和work-item定義為2D
因為work-group和work-item可至3維，get_group_id(), get_global_id(), get_local_id()每一個的參數可以是0~2。
注意，空間維度指數和每個work-group中work-item的數目能夠依據設備而變化。最大維度指數可以通過clGetDeviceInfo()來獲取CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS，每個work-group中work-items的最大值可以通過CL_DEVICE_MAX_WORK_ITEM_SIZES獲取。前者是cl_uint型，后者是size_t的數組。
在kernel函數中，我們能夠通過API調用得到global id以及其他信息：
get_global_id(dim)
get_global_size(dim)
這兩個函數能得到每個維度上的global id。
get_group_id(dim)
get_num_groups(dim)
get_local_id(dim)
get_local_size(dim）
這幾個函數用來計算group id以及在group內的local id。
get_global_id(0) = column, get_global_id(1) = row
get_num_groups(0) * get_local_size(0) == get_global_size(0)
CL_DEVICE_MAX_WORK_ITEM_SIZES，CL_DEVICE_MAX_WORK_GROUP_SIZE（clGetDeivceInfo獲取)它跟CL_KERNEL_WORK_GROUP_SIZE（clGetKernelWorkGroupInfo獲取）有什么區別？
CL_DEVICE_MAX_WORK_ITEM_SIZES ： Max work-items sizes in each dimensions， 每一個維度允許的最大的work-item數
CL_DEVICE_MAX_WORK_GROUP_SIZE： Max allowed work-items in a group,一個workgroup所允許的最多work-item數。
CL_KERNEL_WORK_GROUP_SIZE： Group size returned by kernel 實際在kernel中執行的workgroup數目。
執行cinfo,可以檢測硬件信息