image_scaling.cu
#include <stdio.h>
#include <algorithm>
#include <exception>
#include <iostream>
#include <limits>
#include <sstream>
#include <stdexcept>
#include "cuda_runtime_api.h"
#include "opencv2/opencv.hpp"
#define CHECK(status, lineNum) \
do \
{ \
auto ret = (status); \
if (ret != 0) \
{ \
std::cerr << "Cuda failure: " << cudaGetErrorString(ret) << " lineNum:" << lineNum \
<< std::endl; \
} \
} while (0)
// Kernel which calculate the resized image
__global__ void createResizedImage(unsigned char *imageScaledData,
int scaled_width,
float scale_factor,
cudaTextureObject_t texObj)
{
const unsigned int tidX = blockIdx.x * blockDim.x + threadIdx.x;
const unsigned int tidY = blockIdx.y * blockDim.y + threadIdx.y;
const unsigned index = tidY * scaled_width + tidX;
// Step 5: Read the texture memory from your texture reference in CUDA Kernel
imageScaledData[index] = tex2D<unsigned char>(texObj, (float) (tidX * scale_factor), (float) (tidY * scale_factor));
}
int main(int argc, char *argv[])
{
// Define the scaling ratio
float ResizeScale = 2;
unsigned char *d_scaled_data;
cv::Mat img = cv::imread("../src.jpg", cv::IMREAD_GRAYSCALE);
// cv::resize(img,img, cv::Size(1920, 1024));
int width = img.cols;
int height = img.rows;
std::cout << "src width:" << width << " height:" << height << " img.type():" << img.type()
<< std::endl;
int scaled_height = (int) (height * ResizeScale);
int scaled_width = (int) (width * ResizeScale);
cv::Mat dst = cv::Mat(scaled_height, scaled_width, img.type());
std::cout << "dst width:" << dst.cols << " height:" << dst.rows << " img.type():" << dst.type()
<< std::endl;
// Step 1. Create a channel Description to be used while linking to the tecture
cudaChannelFormatKind kind = cudaChannelFormatKindUnsigned;
/* Returns a channel descriptor using the specified format. */
int x = 8, y = 0, z = 0, w = 0;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(x, y, z, w, kind);
cudaArray *cu_array = nullptr;
CHECK(cudaMallocArray(&cu_array, &channelDesc, width, height), __LINE__);
CHECK(cudaMemcpyToArray(cu_array, 0, 0, img.data, img.total(), cudaMemcpyHostToDevice),
__LINE__);
// Step 2. Specify texture
struct cudaResourceDesc resDesc; //资源描述符,用来获取述纹理数据;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cu_array;
// Step 3. Specify texture object parameters
struct cudaTextureDesc texDesc; //纹理描述符,用来描述纹理参数
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModePoint;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = 0;
// Step 4: Create texture object
cudaTextureObject_t texObj = 0; //需要生产的纹理对象
CHECK(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), __LINE__);
CHECK(cudaMalloc(&d_scaled_data, scaled_height * scaled_width * sizeof(unsigned char)),
__LINE__);
dim3 dimBlock(16, 16, 1);
dim3 dimGrid(scaled_width / dimBlock.x, scaled_height / dimBlock.y, 1);
std::cout << "createResizedImage Launching grid with blocks dimGrid.x:" << dimGrid.x << " dimGrid.y:" << dimGrid.y << std::endl;
createResizedImage<<<dimGrid, dimBlock>>>(d_scaled_data, scaled_width, 1 / ResizeScale, texObj);
CHECK(cudaDeviceSynchronize(), __LINE__);
CHECK(cudaMemcpy(dst.data, d_scaled_data, scaled_height * scaled_width * sizeof(unsigned char),
cudaMemcpyDeviceToHost),
__LINE__);
// Step :6 Destroy texture object
cudaDestroyTextureObject(texObj);
cv::imwrite("dst.jpg", dst);
if (cu_array != NULL)
{
cudaFreeArray(cu_array);
}
if (d_scaled_data != NULL)
{
cudaFree(d_scaled_data);
}
return 0;
}
CMakeLists.txt
cmake_minimum_required(VERSION 3.10)
set(CMAKE_CUDA_HOST_COMPILER " /usr/bin/c++")
# NVCC compiler options
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -ccbin /usr/bin/c++")
project(cudaTest LANGUAGES CUDA CXX)
set(OPENCV_INCLUDE_DIR ./3rdparty/opencv/include CACHE INTERNAL "OPENCV_INCLUDE_DIR")
set(OPENCV_LIB_DIR ./3rdparty/opencv/lib CACHE INTERNAL "OPENCV_LIB_DIR")
if(COMMAND cmake_policy)
cmake_policy(SET CMP0003 NEW)
endif(COMMAND cmake_policy)
add_definitions(-D_FORCE_INLINES)
include_directories(${OPENCV_INCLUDE_DIR})
link_directories(${OPENCV_LIB_DIR})
find_package(CUDA REQUIRED)
if(${CUDA_FOUND})
include_directories(${CUDA_INCLUDE_DIRS})
link_directories($ENV{CUDA_PATH}/lib/x64)
else(${CUDA_FOUND})
MESSAGE(STATUS "cuda not found!")
endif(${CUDA_FOUND})
cuda_add_executable(${PROJECT_NAME}
image_scaling.cu
)
target_link_libraries(
${PROJECT_NAME}
${OPENCV_LIB_DIR}/libopencv_world.so.4.4.0
)