OpenCL與CNN篇二：OpenCL基礎API介紹

阿新 • • 發佈：2019-01-15

本篇介紹幾個OpenCL基礎API，涉及平臺、裝置等初始化必備函式
其次介紹幾個關於緩衝區操作以及工作空間劃分的API
建議閱讀參考書籍，我的推薦是《OpenCL程式設計指南》和《OpenCL異構計算》，尤其是後者從實踐出發，更是適合上手。
本章節函式使用舉例見系列篇三，這裡不再重複；僅介紹API相關引數，以及自己的見解。

一、OpenCL平臺API

平臺API定義了宿主機程式發現OpenCL裝置所用的函式以及這些函式的功能，另外還定義了為OpenCL
應用建立上下文的函式。包括平臺、裝置、上下文等相關函式。
不過為了方便起見這裡將構建程式和建立核心也放在了此章節，即1.4、1.5、1.6。
注意1.7及之後為一些篇外話，不作為和1.1~1.6的連續內容。

1.1、獲取平臺clGetPlatformIDs()

cl_int clGetPlatformIDs( 
       cl_uint          num_entries,    //限制返回的平臺數
       cl_platform_id  *platforms,      //平臺列表儲存位置指標
       cl_uint         *num_platforms   //平臺個數儲存位置指標
        );

將num_entries和platforms分別設定為0和NULL可以查詢可用的平臺個數，儲存在num_platforms中。返回的平臺數可以用num_entries來限制，獲取平臺列表時這個引數要大於0並小於等於可用平臺數。
平臺這個概念應該是不同廠商的實現，比如我一開始安裝了AMD APP SDK，執行程式只有1個可用平臺；後來有想法就又安裝了Intel的SDK，這時執行OpenCL有兩個可用平臺。區別是一個平臺包含GPU和CPU，一個平臺只包含CPU；而且列印裝置資訊，他們包含的CPU名稱相同。

1.2、獲取裝置clGetDeviceIDs()

cl_int clGetDeviceIDs(
       cl_platform_id  platform,       // 指定平臺
       cl_device_type  device_type,    // 指定裝置型別
       cl_uint         num_entries, 
       cl_device_id   *devices,    
       cl_uint        *num_devices
        );

這個命令會得到與platform關聯的eOpenCL裝置列表。如果引數devices為NULL，則用num_devices得到可用裝置數。類似平臺函式用num_entries可以限制返回的裝置數。
引數device_type用來指定計算裝置型別，其極常見引數選擇見下表。

cl_device_type	描述
CL_DEVICE_TYPE_CPU	作為宿主機處理器的OpenCL裝置
CL_DEVICE_TYPE_GPU	作為GPU的OpenCL裝置
CL_DEVICE_TYPE_ACCELERATOR	OpenCL加速器(例如,IBM Cell Broaband)
CL_DEVICE_TYPE_DEFAULT	預設裝置
CL_DEVICE_TYPE_ALL	與相應平臺關聯的所有OpenCL裝置

1.3、建立上下文clCreateContext()

cl_context clCreateContext(
     const cl_context_properties *properties,    // 此處包含平臺資訊
           cl_uint                num_devices,   // 裝置個數
     const cl_device_id          *devices,       // 裝置列表
     void (CL_CALLBACK  *pfn_notify)(const char *errinfo,const void *private_info,size_t cb, void *user_data),
     void                        *user_data,
     cl_int                      *errcode_ret
     );

第一引數著重說明一下，需要用確定的平臺來構建，方法如下：
　　cl_context_properties props[3] =
　　{ CL_CONTEXT_PLATFORM , (cl_context_properties)platformIds, 0 };
第四個引數那麼複雜就直接忽略吧，因為呼叫時那個位置填NULL。使用中user_data也輸入NULL。最後一個引數errcode_ret用來記錄錯誤程式碼，正常時errcode_ret==CL_SUCCESS。

1.4、建立程式物件clCreateProgramWithSource()

cl_program clCreateProgramWithSource(
           cl_context     context,
           cl_uint        count,
           const char   **string,
           const size_t  *lengths,
           cl_int        *errcode_ret
           );

context 建立程式物件的上下文
count 沒看到過介紹
const strings 將這個引數的所有字元，構成了建立程式物件的完整原始碼
lengths 這個引數可以設定為NULL，這種情況下，則認為字串以null終止的
errcode_ret 如果非NULL，函式返回錯誤程式碼將由這個引數返回
一般執行時，先從.cl檔案中讀取程式碼，然後將儲存程式碼的字串傳遞給strings

1.5、構建clBuildProgram()

cl_int clBuildProgram(
       cl_program          program,       // 一個合法的程式物件
       cl_uint             num_devices,   // 要構建程式物件的裝置數
       const cl_device_id *device_list,   // 如果為空則為context所有關聯裝置建立
       const char         *options,
       void(CL_CALLBACK   *pfn_notify)(cl_program program,void *user_data),
       void               *user_data
       );

函式的第四、五、六個引數設定為NULL即可，前三個引數根據程式碼段註釋傳遞即可。深入學習的話關於第四個引數options，可以參考《OpenCL程式設計指南》6.2.2或者檢視文章開頭給的參考網頁。

1.6、建立核心clCreateKernel()

cl_kernel clCreateKernel(
          cl_program  program,
          const char *kernel_name,
          cl_int     *errcode_ret
          );

到這裡就比較容易了，只介紹一下kernel_name：建立核心物件的核心函式名。這是程式原始碼__kernel關鍵字後面的核心函式名。引數errcode_ret的作用和上面的API中意義相同。

1.7、篇外話之先建立上下文再獲取裝置

a、建立上下文clCreateContextFromType()
b、從上下文中查詢裝置資訊clGetContextInfo()

cl_context clCreateContextFromType(
     const cl_context_properties *properties,
     cl_device_type               device_type,
     void (CL_CALLBACK *pfn_notify)(const char *,const void *p,size_t, void *),
     void                        *user_data,
     cl_int                      *errcode_ret           
     );
cl_int clCreatContextInfo(
       cl_Context      context,
       cl_Context_info param_name,
       size_t          param_value_size,
       void           *param_value,
       size_t         *param_value_size_ret
       )

和clCreateContext()類似，不過這裡不需要事先指定裝置列表，使用例項：
　　platform = platformIds[0];
　　cl_context_properties props[3] =
　　{ CL_CONTEXT_PLATFORM,(cl_context_properties)platform, 0 };
　　context = clCreateContextFromType( props, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
之後需要在上下文中查詢裝置clGetContextInfo()，獲取到裝置列表之後操作就和之前相同了，使用例項：
　　cl_uint numDevices ;
　　cl_device_id *deviceIds;
　　errNum = clGetContextInfo( context, CL_CONTEXT_DEVICES, 0, NULL, &numDevices);
　　deviceIds =(cl_device_id*)malloc(sizeof(cl_device_id)*numDevices);
　　errNum = clGetContextInfo( context, CL_CONTEXT_DEVICES, numDevices, deviceIds, NULL);
把這一個篇外話，是因為這有一個不可取之處，這樣搞會有很奇怪的問題。首先在我的電腦上如果是用clCreateContextFromType建立上下文裝置選擇CL_DEVICE_TYPE_CPU；在呼叫clGetContextInfo獲取的numDevices=4，這是我CPU的執行緒數，並不是CPU個數。這樣本來我還能接受，因為只選擇deviceIds[0]的話程式可以正常執行；但如果列印裝置資訊，deviceIds[0]的資訊能正常獲取並且printf出來，但是獲取deviceIds[1]時就開始報錯了，簡直。。。。。
把這一個篇外話，只是為了在實踐中理解上下文和平臺裝置之間的關係，並不是推薦使用。

二、OpenCL執行時API

這些API管理上下文來建立命令佇列以及執行時發生的其他操作。包括建立讀寫緩衝區、設定執行核心等

2.1、設定核心引數clSetKernelArg()

cl_int clSetKernelArg(
       cl_kernel   kernel,
       cl_uint     arg_index,
       size_t     *arg_size,
       const void *arg_value
       );

kernel 一個合法的核心物件；arg_index 核心引數的索引；arg_size 引數的大小；arg_value 傳入核心函式的引數的一個指標。這裡引數的意思看看就好，重要的是看例子學會使用。設定好kernel引數之後通過呼叫下文的clEnqueueNDRangeKernel來執行。
值得說明的是如果你要多次呼叫執行同一個核心，那麼可以只設置一次核心。例如我寫卷積神經網路，在程式開始設定核心引數，之後訓練過程中只需把新的影象寫入到影象緩衝區，直接呼叫clEnqueueNDRangeKernel執行核心就好。

2.2、執行核心clEnqueueNDRangeKernel()

cl_int clEnqueueNDRangeKernel(
       cl_command_queue commad_queue,
       cl_kernel        kernel,
       cl_uint          work_dim,
       const size_t    *global_work_offset,
       const size_t    *global_work_size,
       const size_t    *local_work_size,
       cl_uint          num_events_in_wait_list,
       const cl_event  *events_in_wait_list,
       cl_event        *event
       );

commend_queue 核心的執行需要提交到命令佇列中
kernel 要執行的核心名稱
work_dim 指定新建的work-item的緯度，這裡我假設你已經知道了OpenCL的核心與工作項之間的關係。
global_work_offset 為work-item提供全域性ID，該引數可以不從0開始，但一般設定為0或者NULL
global_work_size 指定NDRange中每維work-item的數量，不可為空
local_work_size 指定workgroup中每維work-item的數量，可以設定為NULL讓系統自動設定
num_events_in_wait_list、events_in_wait_list、event 這是OpenCL高階一點的操作，用於記錄此事件或者需要等待的其他事件，可以用來規劃不同事件執行順序；如果不用可以分別設定為0、NULL、NULL。
需要注意的是global_work_size指向的陣列大小要和work_dim相等，global_work_offset和local_work_size不為NULL時也是同樣要求，不一樣會發生什麼我也沒有測試，不過正常的程式這裡該是一樣的。

三、緩衝區的建立、寫入和讀取

3.1、flush和finish命令

這兩個命令在命令佇列中值兩種不同型別的barrier操作。clFinish()函式阻塞直到命令佇列中所有命令完成。clFlush()阻塞指導命令佇列中的所有命令被移出佇列，這意味著這些命令已經準備就緒但無法保證執行完畢。
cl_int clFlush(cl_command_queue command_queue);
cl_int clFinish(cl_command_queue command_queue);

3.2、建立緩衝區clCreateBuffer()

cl_mem clCreateBuffer(
       cl_context   context,
       cl_mem_flags flag,
       size_t       size,
       void        *host_ptr,
       cl_int      *errcode_ret
       );

context 一個合法的上下文，為這個上下文分配緩衝區
size 所分配緩衝區的大小（位元組數）
host_ptr 這個指標在clCreateBuffer如何使用有flags引數確定。host_ptr指向的資料大小應大於等於size
flags 用於指定關於緩衝區建立的分配和使用資訊。其部分取值見下表。

cl_mem_flags	描述
CL_MEM_READ_WRITE	指定記憶體物件將由核心讀寫；預設為此模式
CL_MEM_WRITE_ONLY	指定記憶體物件由核心寫，但不能讀。
CL_MEM_READ_ONLY	指定記憶體物件由核心讀，但不能寫。
CL_MEM_USE_HOST_PTR	只有當host_ptr為非NULL時，這個標誌合法；使用host_ptr引用的記憶體作為記憶體物件的儲存位
CL_MEM_ALLOC_HOST_PTR	指定緩衝區應當在宿主機可訪問的記憶體中分配。不可與USE_HOST_PTR同時使用
CL_MEM_COPY_HOST_PTR	表示希望OpenCL實現分配記憶體物件的記憶體，並從host_ptr引用複製資料。

3.3、讀寫緩衝區

cl_int clEnqueueWriteBuffer(
       cl_command_queue commad_queue,
       cl_mem           buffer,
       cl_bool          blocking_write,
       size_t           offset,
       size_t           cb,
       void            *ptr,
       cl_uint          num_events_in_wait_list,
       const cl_event  *events_in_wait_list,
       cl_event        *event
       );
cl_int clEnqueueReadBuffer(
       cl_command_queue commad_queue,
       cl_mem           buffer,
       cl_bool          blocking_read,
       size_t           offset,
       size_t           cb,
       void            *ptr,
       cl_uint          num_events_in_wait_list,
       const cl_event  *events_in_wait_list,
       cl_event        *event
          );

command_queue 這是一個命令佇列,讀寫命令將在這個佇列中排隊
buffer 一個合法的緩衝區物件（資料對這裡讀寫）
blocking_read 如果設定為CL_TRUE，則命令阻塞,直至ptr讀寫資料完成
offset 緩衝區物件的讀寫資料的起始偏移量（位元組數）
cb 對緩衝區讀寫的位元組數
ptr 宿主機記憶體的一個指標，寫入緩衝區的資料從哪裡來 / 或者從緩衝區讀資料寫入哪裡
關於讀寫的blocking_write的多一些說明：
　　如果blocking_write為CL_TRUE,
　　　　則OpenCL實現將複製ptr引用的資料，並在命令佇列中對寫操作進行排隊。
　　　　在clEnqueueWriteBuffer呼叫返回後，由ptr指向的記憶體可以被應用程式重用。
　　如果blocking_write為CL_FALSE,
　　　　則OpenCL實現將使用ptr執行非阻塞寫操作。由於寫是非阻塞的，實現可以立即返回。
　　　　ptr指向的記憶體在呼叫返回後不能被應用程式重用。
　　　　event引數返回一個事件物件，可以用來查詢write命令的執行狀態。
　　　　當寫命令完成後，ptr指向的記憶體可以被應用程式重新使用

四、影象建立讀寫

影象型別的資料我還沒有用過，不過好像挺有用的；在此只給出API的介紹。

4.1、影象格式

typedef struct _cl_image_format
{
    cl_channel_order image_channel_order;
    cl_channel_type  iamge_channel_date_type; 
}   cl_image_format ;

The image format describes how the data will be stored in memory
使用示例:
cl_image_format format;
format.image_channel_order = CL_R; // single channel
format.image_channel_data_type = CL_FLOAT; // float data type

4.2、影象建立

cl_mem clCreateImage2D(
       cl_context             context,
       cl_mem_flags           flags,
       const cl_image_format *image_format
       size_t                 image_with,
       size_t                 image_height,
       size_t                 image_row_pitch,
       void                  *host_ptr,
       cl_int                *errcode_ret
          );
cl_mem clCreateImage3D(
       cl_context             context,
       cl_mem_flags           flags,
       const cl_image_format *image_format
       size_t                 image_with,
       size_t                 image_height,
       size_t                 image_depth,
       size_t                 image_row_pitch,
       size_t                *image_slice_pitch,
       void                  *host_ptr,
       cl_int                *errcode_ret
          );

context 建立影象物件的上下文
flags 其合法列舉由cl_mem_flags定義
image_format 描述通道次序和影象通道資料型別
image_with,image_height,image_depth 影象的長寬深
image_row_pitch 如果host_ptr不為NULL,這個值指定影象中各行的位元組數;為0採取預設值
image_slice_pitch 如果host_ptr不為NULL,這個值指定影象中各個切片的位元組數;為0採取預設值
host_ptr 記憶體中線性佈局的影象緩衝區指標
errcode_ret 如果為非NULL,函式返回錯誤碼
使用示例:
cl_mem d_inputImage = clCreateImage2D(context, 0, &format, imageWidth, imageHeight, 0, NULL, &errNum);

4.3、影象讀寫

cl_int clEnqueueWriteImage(
       cl_command_queue commad_queue,
       cl_mem           image,
       cl_bool          blocking_read,
       const size_t     origin[3],
       const size_t     region[3],
       size_t           row_pitch
       size_t           slice_pitch,
       void            *ptr,
       cl_uint          num_events_in_wait_list
       const cl_event  *event_wait_list,
       cl_event        *event
          );

commad_queue 寫入命令將放入這個佇列
iamge 這是一個合法的影象物件
blocking_read 如果設定為CL_TRUE,則clEnqueueReadImage阻塞,直到資料讀入ptr
origin 要寫入相對影象原點的(x,y,z)整數座標,對於二維z=0
region 要寫入區域的(寬,高,深),對於二維z=1
row_pitch 影象中各行位元組數,預設為image_with*(byte_per_pixel)
slice_pitch 三維影象中各切片的位元組數image_height*row_pitch
ptr 這個指標指向源資料的宿主機記憶體
使用示例:
errNum = clEnqueueWriteImage(queue, d_inputImage, CL_FALSE, origin, region, 0, 0, inputImage, 0, NULL, NULL);

cl_int clEnqueueReadImage(
       cl_command_queue commad_queue,
       cl_mem           image,
       cl_bool          blocking_read,
       const size_t     origin[3],
       const size_t     region[3],
       size_t           row_pitch
       size_t           slice_pitch,
       void            *ptr,
       cl_uint          num_events_in_wait_list
       const cl_event  *event_wait_list,
       cl_event        *event
        );

commad_queue 讀取命令將放入這個佇列
iamge 這是一個合法的影象物件
blocking_read 如果設定為CL_TRUE,則clEnqueueReadImage阻塞,直到資料讀入ptr
origin 要讀取相對原影象原點的(x,y,z)整數座標,對於二維z=0
region 要讀取區域的(寬,高,深),對於二維z=1
row_pitch 影象中各行位元組數,預設為image_with*(byte_per_pixel)
slice_pitch 三維影象中各切片的位元組數image_height*row_pitch
ptr 這個指標指向寫入資料的宿主機記憶體
使用示例:
errNum = clEnqueueReadImage(queue, d_outputImage, CL_TRUE, origin, region, 0, 0, outputImage, 0, NULL, NULL);

4.4、篇外話

1.要實現影象類資料的使用，好像還需要宣告一個取樣器（cl_sampler），具體怎麼操作我還沒有試驗過。
2.在核心中讀寫影象好像需要使用固定的函式read_imagef()和write_imagef()；具體解釋參考OpenCL Reference Pages =>OpenCL Compiler =>Built-in Functions => Image Functions =>read_imagef / write_imagef 。