- 操作系統：ubuntu22.04
- OpenCV版本：OpenCV4.9
- IDE:Visual Studio Code
- 編程語言：C++11
算法描述
在同一個線程塊（thread block內，將 [beg, end) 范圍內的數據并行地復制到 out 開始的位置。
它使用了 CUDA 線程協作機制（warp-level 或 block-level） 來實現高效的塊級拷貝，通常比簡單的逐線程拷貝更快。
函數原型
_device__ static __forceinline__ void cv::cudev::blockCopy
(InIt beg,InIt end,OutIt out
)
參數
| 參數名 | 類型 | 含義 |
|---|---|---|
| beg | InIt | 輸入范圍的起始迭代器（或指針） |
| end | InIt | 輸入范圍的結束迭代器（不包含該位置） |
| out | OutIt | 輸出范圍的起始迭代器（或指針） |
使用場景
- 圖像處理中的 局部塊拷貝
- 構建 GPU 并行歸約（reduction）算法前的數據準備
- 實現滑動窗口（sliding window）或卷積操作時的 tile 數據加載
- 需要多個線程協同完成連續內存拷貝的任務
示例代碼
#include <opencv2/cudev/functional/functional.hpp>
#include <cstdio>
#include <opencv2/core.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/cudev/block/block.hpp>
#include <iostream>using namespace cv::cudev;// 定義 tile 大小為 16x16
#define TILE_SIZE 16__global__ void processImageKernel(const uchar* input, size_t pitchIn,uchar* output, size_t pitchOut,int width, int height) {// 共享內存用于存儲當前線程塊處理的 tile__shared__ uchar tile[TILE_SIZE][TILE_SIZE];// 當前線程負責的圖像坐標int x = blockIdx.x * TILE_SIZE + threadIdx.x;int y = blockIdx.y * TILE_SIZE + threadIdx.y;// 檢查是否在圖像范圍內if (x >= width || y >= height)return;// 輸入和輸出指針const uchar* in_ptr = input + y * pitchIn + x;uchar* out_ptr = output + y * pitchOut + x;// 將當前 tile 拷貝到共享內存blockCopy(in_ptr, in_ptr + TILE_SIZE * TILE_SIZE, &tile[0][0]);__syncthreads(); // 必須同步所有線程，確保共享內存拷貝完成// 對 tile 中的數據進行簡單處理（比如增加亮度）uchar value = tile[threadIdx.y][threadIdx.x];value = min(value + 50, (uchar)255);// 寫回到輸出位置blockCopy(&value, &value + 1, out_ptr);
}int main() {// 創建一個測試圖像（8UC1）cv::Mat h_src = cv::Mat::zeros(128, 128, CV_8UC1);h_src.setTo(100); // 填充為灰度值 100// 上傳到 GPUcv::cuda::GpuMat d_src, d_dst;d_src.upload(h_src);d_dst.create(h_src.size(), h_src.type());// 設置 kernel 參數dim3 threads(TILE_SIZE, TILE_SIZE);dim3 blocks((h_src.cols + TILE_SIZE - 1) / TILE_SIZE,(h_src.rows + TILE_SIZE - 1) / TILE_SIZE);// 調用 kernelprocessImageKernel<<<blocks, threads>>>(d_src.ptr<uchar>(), d_src.step,d_dst.ptr<uchar>(), d_dst.step,d_src.cols, d_dst.rows);// 下載結果cv::Mat h_dst;d_dst.download(h_dst);// 顯示部分結果std::cout << "Output image sample:\n" << h_dst(cv::Rect(0, 0, 5, 5)) << std::endl;return 0;
}
運行結果
Output image sample:
[150, 0, 0, 0, 0;0, 0, 0, 0, 0;0, 0, 0, 0, 0;0, 0, 0, 0, 0;0, 0, 0, 0, 0]
