Đây là trường hợp cổ điển của embarrassingly parallel vấn đề xử lý hình ảnh có thể dễ dàng được ánh xạ tới khung công tác CUDA. Bộ lọc trung bình được biết là Box Filter trong các miền xử lý hình ảnh.
Cách tiếp cận dễ nhất là sử dụng kết cấu CUDA cho quy trình lọc vì điều kiện biên có thể được xử lý rất dễ dàng bằng kết cấu.
Giả sử bạn có con trỏ nguồn và đích được phân bổ trên máy chủ lưu trữ. Các thủ tục sẽ là một cái gì đó như thế này.
- Phân bổ bộ nhớ đủ lớn để giữ hình ảnh nguồn và đích trên thiết bị.
- Sao chép hình ảnh nguồn từ máy chủ sang thiết bị.
- Liên kết con trỏ thiết bị nguồn hình ảnh với kết cấu.
- Chỉ định kích thước khối thích hợp và lưới đủ lớn để bao phủ mọi pixel của hình ảnh.
- Khởi chạy hạt nhân lọc bằng cách sử dụng lưới và kích thước khối được chỉ định.
- Sao chép kết quả về máy chủ.
- Bỏ kết cấu
- Con trỏ thiết bị miễn phí.
mẫu thực hiện của hộp lọc
Kernel
texture<unsigned char, cudaTextureType2D> tex8u;
//Box Filter Kernel For Gray scale image with 8bit depth
__global__ void box_filter_kernel_8u_c1(unsigned char* output,const int width, const int height, const size_t pitch, const int fWidth, const int fHeight)
{
int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
const int filter_offset_x = fWidth/2;
const int filter_offset_y = fHeight/2;
float output_value = 0.0f;
//Make sure the current thread is inside the image bounds
if(xIndex<width && yIndex<height)
{
//Sum the window pixels
for(int i= -filter_offset_x; i<=filter_offset_x; i++)
{
for(int j=-filter_offset_y; j<=filter_offset_y; j++)
{
//No need to worry about Out-Of-Range access. tex2D automatically handles it.
output_value += tex2D(tex8u,xIndex + i,yIndex + j);
}
}
//Average the output value
output_value /= (fWidth * fHeight);
//Write the averaged value to the output.
//Transform 2D index to 1D index, because image is actually in linear memory
int index = yIndex * pitch + xIndex;
output[index] = static_cast<unsigned char>(output_value);
}
}
Wrapper Chức năng:
void box_filter_8u_c1(unsigned char* CPUinput, unsigned char* CPUoutput, const int width, const int height, const int widthStep, const int filterWidth, const int filterHeight)
{
/*
* 2D memory is allocated as strided linear memory on GPU.
* The terminologies "Pitch", "WidthStep", and "Stride" are exactly the same thing.
* It is the size of a row in bytes.
* It is not necessary that width = widthStep.
* Total bytes occupied by the image = widthStep x height.
*/
//Declare GPU pointer
unsigned char *GPU_input, *GPU_output;
//Allocate 2D memory on GPU. Also known as Pitch Linear Memory
size_t gpu_image_pitch = 0;
cudaMallocPitch<unsigned char>(&GPU_input,&gpu_image_pitch,width,height);
cudaMallocPitch<unsigned char>(&GPU_output,&gpu_image_pitch,width,height);
//Copy data from host to device.
cudaMemcpy2D(GPU_input,gpu_image_pitch,CPUinput,widthStep,width,height,cudaMemcpyHostToDevice);
//Bind the image to the texture. Now the kernel will read the input image through the texture cache.
//Use tex2D function to read the image
cudaBindTexture2D(NULL,tex8u,GPU_input,width,height,gpu_image_pitch);
/*
* Set the behavior of tex2D for out-of-range image reads.
* cudaAddressModeBorder = Read Zero
* cudaAddressModeClamp = Read the nearest border pixel
* We can skip this step. The default mode is Clamp.
*/
tex8u.addressMode[0] = tex8u.addressMode[1] = cudaAddressModeBorder;
/*
* Specify a block size. 256 threads per block are sufficient.
* It can be increased, but keep in mind the limitations of the GPU.
* Older GPUs allow maximum 512 threads per block.
* Current GPUs allow maximum 1024 threads per block
*/
dim3 block_size(16,16);
/*
* Specify the grid size for the GPU.
* Make it generalized, so that the size of grid changes according to the input image size
*/
dim3 grid_size;
grid_size.x = (width + block_size.x - 1)/block_size.x; /*< Greater than or equal to image width */
grid_size.y = (height + block_size.y - 1)/block_size.y; /*< Greater than or equal to image height */
//Launch the kernel
box_filter_kernel_8u_c1<<<grid_size,block_size>>>(GPU_output,width,height,gpu_image_pitch,filterWidth,filterHeight);
//Copy the results back to CPU
cudaMemcpy2D(CPUoutput,widthStep,GPU_output,gpu_image_pitch,width,height,cudaMemcpyDeviceToHost);
//Release the texture
cudaUnbindTexture(tex8u);
//Free GPU memory
cudaFree(GPU_input);
cudaFree(GPU_output);
}
Tin tốt là bạn không cần phải thực hiện tự lọc. Bộ công cụ CUDA đi kèm với thư viện xử lý hình ảnh và tín hiệu miễn phí có tên là NVIDIA Performance Primitives, còn gọi là NPP, do NVIDIA sản xuất. NPP sử dụng CUDA cho phép GPU tăng tốc quá trình xử lý. Bộ lọc trung bình đã được triển khai trong NPP. Phiên bản hiện tại của NPP (5.0) hỗ trợ hình ảnh 8 bit, 1 kênh và 4 kênh. Các chức năng là:
nppiFilterBox_8u_C1R cho hình ảnh 1 kênh.nppiFilterBox_8u_C4R cho hình ảnh 4 kênh.
Xin chào @SamehKamal, tôi tò mò tò mò. Các mã sử dụng CUDA nhanh như thế nào so với các mã truyền thống trong kết quả của bạn? –
Đây là một thời đại dài tôi không nhớ chính xác yếu tố tăng tốc cho thuật toán này nhưng hiệu suất thay đổi từ thuật toán này sang thuật toán khác từ tỷ lệ tăng tốc x7 đến x22 cho thuật toán tôi đã sử dụng. –