OpenCL：影象處理基礎note
使用影象物件的理由
雖然對於影象也可以把它的畫素資料當做一般的快取資料來處理，但是如果把它當做影象來處理有如下好處：
- 在GPU中，影象資料是儲存在特殊的全域性記憶體中，即紋理記憶體，它和一般的全域性記憶體不相同，它是被快取的，用於高速訪問處理。GPU中有專門支援影象讀寫的硬體，使用內建讀寫函式可以充分發揮這個優勢。
- 只要OpenCL支援該影象格式，那麼就可以不用考慮影象格式的前提下使用讀寫影象資料的函式
- 可以使用取樣器來配置讀取影象中資料的的方式
- OpenCL提供函式來獲取影象相關資訊，比如寬度等
主機與核心的命名
影象相關應用主要涉及兩個資料型別，即影象物件和取樣器。影象對線用來儲存主機和裝置上的影象物件，而取樣器則在裝置接收資料時，說明如何讀取這些顏色值
| 影象物件 | 取樣器 | |
|---|---|---|
| 主機 | cl_mem (和一般快取物件一樣) | cl_sampler |
| 裝置 | image2d_t或者image3d_t | sampler_t |
主機程式設計的主要介面
建立影象物件
cl_mem clCreateImage2D()
cl_mem clCreateImage3D()
影象格式
cl_image_format
建立取樣器
cl_sampler clCreateSampler()
裝置上的介面
影象資料
由於很多裝置都是將影象物件儲存在特定的記憶體中，因此image2d_t以及image3d_t資料的前面通常會有read_only或者write_only等修飾符。影象不能即可寫又可讀
影象資料作為核心的引數時，和普通的快取物件不同，它不是指標引數，因為影象資料並不希望直接通過記憶體操作來訪問資料。
取樣器
根據khronos官網的說明，取樣器其實就是一個uint型別（可以直接給內建cl讀寫函式傳遞標識位的組合來替代sampler_t，例如：float4 clr = read_imagef(img, CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST, (float2)(x, y));），通過bit-mode來說明座標模式，超出邊界的處理模式以及插值模式。在官方的sample中採用的是constant申明的模式，雖然也說了可以採用global宣告，但這也許說明constant更好。
可以通過設定引數clSetKernelArg
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
__kernel void simple_image(read_only image2d_t src_image,
write_only image2d_t dst_image) {
/* Compute value to be subtracted from each pixel */
uint offset = get_global_id(1) * 0x4000 + get_global_id(0) * 0x1000;
/* Read pixel value */
int2 coord = (int2)(get_global_id(0), get_global_id(1));
uint4 pixel = read_imageui(src_image, sampler, coord);
/* Subtract offset from pixel */
pixel.x -= offset;
/* Write new pixel value to output */
write_imageui(dst_image, coord, pixel);
}
內建影象函式
OpenCL的內建影象相關的函式主要包括三類：read_imageT讀取類，write_imageT寫入類以及get_image_X讀取影象資訊類。其中T表示資料型別，比如f，i，ui等，而X表示width，dim等影象資訊
read_image
- read_image函式返回的都是四分量的資料，如果實際圖片沒有四個通道的資料，那麼那些不含有的的通道對應的分量會被置為0，具體請看官網
- 很多時候都會看到程式碼中用
float4 read_imagef()去讀影象資料為int或者uint型別的image，可能你會困惑為什麼整形要用float去讀？這是因為，這些圖片在建立的時候就是以正則NORM化的資料型別，這個型別在cl_image_format的image_channel_data_type中指定。以int8為例，image_channel_data_type可取的值如下表：參考
| image_channel_data_type 引數 | 含義 | read_image |
|---|---|---|
| CL_SNORM_INT8 | 每個通道都是正則化的有符號8位整數 | 使用read_imagef讀取，返回的float4的每個分量的取值範圍為[0.0, 1.0] |
| CL_UNORM_INT8 | 每個通道都是正則化的無符號8位整數 | 使用read_imagef讀取，返回的float4的每個分量的取值範圍為[-1.0, 1.0] |
| CL_SIGNED_INT8 | 每個通道都是8位有符號的整數 | 使用read_imagei讀取，返回int4 |
| CL_USIGNED_INT8 | 每個通道都是8位無符號的整數 | 使用read_imageui讀取，返回unsigned int4 |
- 使用read_imagei和read_imageui時，取樣器的filter_mode必須設定為CLK_FILTER_NEAREST
- 使用read_imagef時，如果座標是整數，那麼filte_mode必須設定為CLK_FILTER_NEAREST
- 使用read_imageX時，X表示**，如果座標是整數，那麼取樣器的座標正則化必須設定為CLK_NORMALIZED_COORDS_FALSE，且address mode必須設定為CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP 或者 CLK_ADDRESS_NONE
使用影象物件的理由
雖然對於影象也可以把它的畫素資料當做一般的快取資料來處理，但是如果把它當做影象來處理有如下好處：
- 在GPU中，影象資料是儲存在特殊的全域性記憶體中，即紋理記憶體，它和一般的全域性記憶體不相同，它是被快取的，用於高速訪問處理。GPU中有專門支援影象讀寫的硬體，使用內建讀寫函式可以充分發揮這個優勢。
- 只要OpenCL支援該影象格式，那麼就可以不用考慮影象格式的前提下使用讀寫影象資料的函式
- 可以使用取樣器來配置讀取影象中資料的的方式
- OpenCL提供函式來獲取影象相關資訊，比如寬度等
主機與核心的命名
影象相關應用主要涉及兩個資料型別，即影象物件和取樣器。影象對線用來儲存主機和裝置上的影象物件，而取樣器則在裝置接收資料時，說明如何讀取這些顏色值
| 影象物件 | 取樣器 | |
|---|---|---|
| 主機 | cl_mem (和一般快取物件一樣) | cl_sampler |
| 裝置 | image2d_t或者image3d_t | sampler_t |
主機程式設計的主要介面
建立影象物件
cl_mem clCreateImage2D()
cl_mem clCreateImage3D()
影象格式
cl_image_format
建立取樣器
cl_sampler clCreateSampler()
裝置上的介面
影象資料
由於很多裝置都是將影象物件儲存在特定的記憶體中，因此image2d_t以及image3d_t資料的前面通常會有read_only或者write_only等修飾符。影象不能即可寫又可讀。
影象資料作為核心的引數時，和普通的快取物件不同，它不是指標引數，因為影象資料並不希望直接通過記憶體操作來訪問資料。
取樣器
根據khronos官網的說明，取樣器其實就是一個uint型別（可以直接給內建cl讀寫函式傳遞標識位的組合來替代sampler_t，例如：float4 clr = read_imagef(img, CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST, (float2)(x, y));），通過bit-mode來說明座標模式，超出邊界的處理模式以及插值模式。在官方的sample中採用的是constant申明的模式，雖然也說了可以採用global宣告，但這也許說明constant更好。
可以通過設定引數clSetKernelArg的方式來給核心傳遞取樣器。但是更為方便的方式是在cl檔案中，在核心函式之前建立一個sampler_t物件，如下case：
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
__kernel void simple_image(read_only image2d_t src_image,
write_only image2d_t dst_image) {
/* Compute value to be subtracted from each pixel */
uint offset = get_global_id(1) * 0x4000 + get_global_id(0) * 0x1000;
/* Read pixel value */
int2 coord = (int2)(get_global_id(0), get_global_id(1));
uint4 pixel = read_imageui(src_image, sampler, coord);
/* Subtract offset from pixel */
pixel.x -= offset;
/* Write new pixel value to output */
write_imageui(dst_image, coord, pixel);
}
內建影象函式
OpenCL的內建影象相關的函式主要包括三類：read_imageT讀取類，write_imageT寫入類以及get_image_X讀取影象資訊類。其中T表示資料型別，比如f，i，ui等，而X表示width，dim等影象資訊
read_image
- read_image函式返回的都是四分量的資料，如果實際圖片沒有四個通道的資料，那麼那些不含有的的通道對應的分量會被置為0，具體請看官網
- 很多時候都會看到程式碼中用
float4 read_imagef()去讀影象資料為int或者uint型別的image，可能你會困惑為什麼整形要用float去讀？這是因為，這些圖片在建立的時候就是以正則NORM化的資料型別，這個型別在cl_image_format的image_channel_data_type中指定。以int8為例，image_channel_data_type可取的值如下表：參考
| image_channel_data_type 引數 | 含義 | read_image |
|---|---|---|
| CL_SNORM_INT8 | 每個通道都是正則化的有符號8位整數 | 使用read_imagef讀取，返回的float4的每個分量的取值範圍為[0.0, 1.0] |
| CL_UNORM_INT8 | 每個通道都是正則化的無符號8位整數 | 使用read_imagef讀取，返回的float4的每個分量的取值範圍為[-1.0, 1.0] |
| CL_SIGNED_INT8 | 每個通道都是8位有符號的整數 | 使用read_imagei讀取，返回int4 |
| CL_USIGNED_INT8 | 每個通道都是8位無符號的整數 | 使用read_imageui讀取，返回unsigned int4 |
- 使用read_imagei和read_imageui時，取樣器的filter_mode必須設定為CLK_FILTER_NEAREST
- 使用read_imagef時，如果座標是整數，那麼filte_mode必須設定為CLK_FILTER_NEAREST
- 使用read_imageX時，X表示**，如果座標是整數，那麼取樣器的座標正則化必須設定為CLK_NORMALIZED_COORDS_FALSE，且address mode必須設定為CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP 或者 CLK_ADDRESS_NONE