CUDA编程模型基础01

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
)