OpenCL 第10課:kernel,work_item和workgroup

轉載自：http://www.cmnsoft.com/wordpress/?p=1429

前幾節我們一起學習了幾個用ＯＰＥＮＣＬ完成任務的簡單例子，從這節起我們將更詳細的對ＯＰＥＮＣＬ進行一些“理論”學習。

 

kernel：是指一個用opencl c語言編寫的、代表一個單一執行實例的代碼單元。opencl 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在全局空間中的索引。

關於使用這些函數，我們舉一個之前學過的例子。在第７課《旋轉變換（１）》中的內核程序中原文是這樣的。

__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),程序運行結果是一樣的。而且我們還可以少傳遞兩個參數。節省空間，提高效率。大家看下以前的例子，看看那些代碼我們還可以優化呢。