OpenCL 第10课:kernel,work_item和workgroup
轉載自:http://www.cmnsoft.com/wordpress/?p=1429
前幾節我們一起學習了幾個用OPENCL完成任務的簡單例子,從這節起我們將更詳細的對OPENCL進行一些“理論”學習。
kernel:是指一個用opencl c語言編寫的、代表一個單一執行實例的代碼單元。opencl c語言看起來跟C語言函數非常相像,都有一個參數列表“局部”變量定義和標準控制流結構。opencl術語中把這種kernel實例稱為work-item(工作項)。但opencl kernel與c語方函數的區別在于其并行語義。
work_item:是定義在一個很大的并行執行空間中的一小部分。是并行操作中每一部分的實例化。通俗來說,可以理解為kernel里定義的執行函數。當kernel啟動后會創建大量work_item來同時執行,以完成并行任務。work_item根所其數據結構大小可分為一維、二維和三維數據。work_item之是的運行是相互獨立的,不同步的。
work_group:opencl將全局執行空間劃分為大量大小相等的,一維、二維、三維的work_item集合,這個集合就是work_group。在work_group內部,各個work_item之間允許一定程度的通信。而有work_group保證并發執行來允許其內部的work_item之間的本地同步。
在實際編寫內核中,要了解線程調度的維度數,work_group的大小是很重要的,這有利于我們優化編寫的內核程序。opencl提供了一此非常有用的函數供我們調用(在內核中調用)。
uint get_work_dim(): 返回線程調度的維度數。
uint get_global_size(uint dimension): 返回在所請求維度上work_item的總數。
uint get_global_id(uint dimension): 返回在所請求的維度上當前work_item在全局空間中的索引。
uint get_local_size(uint dimension): 返回在所請求的維度上work-group的大小。
uint get_local_id(uint dimension): 返回在所請求的維度上,當前work_item在work_group中的索引。
uint get_number_groups(uint dimension): 返回在所請求維度上work-group的數目,這個值等于get_global_size 除以 get_local_size。
uint get_group_id(uint dimension): 返回在所請求的維度上當前wrok_group在全局空間中的索引。
關于使用這些函數,我們舉一個之前學過的例子。在第7課《旋轉變換(1)》中的內核程序中原文是這樣的。
__kernel void rotation(__global int* A,
__global int* B,
int width,
int height,
float sinangle,
float cosangle)
{
//獲取索引號,這里是二維的,所以可以取兩個
//否則另一個永遠是0
int col = get_global_id(0);
int row = get_global_id(1);
//計算圖形中心點
float cx = ((float)width)/2;
float cy = ((float)height)/2;
int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
int ny = (int)(cy + (-1*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy));
//邊界檢測
if(nx>=0 && nx<width && ny>=0 && ny<height)
{
B[nx + ny*width] = A[col + row*width];
}
}
這里傳遞的width和height大小是一樣的,表示圖像數據長寬的大小。其實也就是維度上work_item的總數,我們可以把代碼改成。
__kernel void rotation(__global int* A,
__global int* B,
int width,
int height,
float sinangle,
float cosangle)
{
//獲取索引號,這里是二維的,所以可以取兩個
//否則另一個永遠是0
int col = get_global_id(0);
int row = get_global_id(1);
//計算圖形中心點
float cx = ((float)get_global_size(0))/2;
float cy = ((float)get_global_size(0))/2;
int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
int ny = (int)(cy + (-1*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy));
//邊界檢測
if(nx>=0 && nx<get_global_size(0) && ny>=0 && ny<get_global_size(0))
{
B[nx + ny*get_global_size(0)] = A[col + row*get_global_size(0)];
}
}
把所有的width和heigh全部改成get_global_size(0),程序運行結果是一樣的。而且我們還可以少傳遞兩個參數。節省空間,提高效率。大家看下以前的例子,看看那些代碼我們還可以優化呢。
總結
以上是生活随笔為你收集整理的OpenCL 第10课:kernel,work_item和workgroup的全部內容,希望文章能夠幫你解決所遇到的問題。
- 上一篇: 大眼橙上架 C1Air 投影仪:“真 1
- 下一篇: 华为松山湖运动科学实验室揭秘:华为穿戴领