0_Simple__simpleSurfaceWrite

使用表面写入函数，结合纹理引用实现图片的旋转
▶ 源代码

#include <stdio.h>
#include <windows.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_functions.h>
#include <helper_cuda.h>    

#define WINDOWS_LEAN_AND_MEAN
#define NOMINMAX
#define MIN_EPSILON_ERROR 5e-3f
float angle = 0.5f;                             // 弧度制
texture<float, 2, cudaReadModeElementType> tex;
surface<void, 2> outputSurface;

// 使用表面写入，将全局内存中的数据 d_data 写到绑定了纹理引用的 CUDA 数组 cuArray 中
__global__ void surfaceWriteKernel(float *gIData, int width, int height)
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    surf2Dwrite(gIData[y * width + x], outputSurface, x * 4, y, cudaBoundaryModeTrap);
}

// 利用纹理取样，将绑定了纹理引用的 CUDA 数组 cuArray 中的图片进行旋转，写入全局内存 d_data 中
__global__ void transformKernel(float *gOData,int width,int height,float theta) 
{
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;    
    float u = x / (float)width - 0.5f;
    float v = y / (float)height - 0.5f;

    gOData[y * width + x] = tex2D(tex, u * cosf(theta) - v * sinf(theta) + 0.5f, v * cosf(theta) + u * sinf(theta) + 0.5f);
}

int main()
{
    printf("
	Start.
");
    cudaSetDevice(0);// 删掉了筛选设备的过程
    cudaDeviceProp deviceProps;
    cudaGetDeviceProperties(&deviceProps, 0);
    printf("
	Device %s, Multi-Processors: %d, SM %d.%d
", deviceProps.name, deviceProps.multiProcessorCount, deviceProps.major, deviceProps.minor);

    // 读取图片数据
    float *h_data = NULL, *h_dataRef = NULL;
    unsigned int width, height, size;
    sdkLoadPGM("D:\Code\CUDA\cudaProjectTemp\data\lena_bw.pgm", &h_data, &width, &height);// 删掉了用函数 sdkFindFilePath() 查找输入文件的过程
    size = width * height * sizeof(float);
    sdkLoadPGM("D:\Code\CUDA\cudaProjectTemp\data\ref_rotated.pgm", &h_dataRef, &width, &height);
    printf("
	Load input files, %d x %d pixels
", width, height);

    // 申请设备内存
    float *d_data = NULL;
    cudaMalloc((void **) &d_data, size);
    cudaArray *cuArray;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaMallocArray(&cuArray,&channelDesc,width,height,cudaArraySurfaceLoadStore);
    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
    //cudaMemcpyToArray(cuArray,0,0,h_data,size,cudaMemcpyHostToDevice); 只使用纹理内存时，可以直接拷贝到cuArray中
    
    // 绑定表面引用
    cudaBindSurfaceToArray(outputSurface, cuArray, channelDesc);

    // 使用表面写入
    dim3 dimBlock(8, 8, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    surfaceWriteKernel<<<dimGrid, dimBlock>>>(d_data, width, height);

    // 绑定纹理引用
    tex.addressMode[0] = cudaAddressModeWrap;
    tex.addressMode[1] = cudaAddressModeWrap;
    tex.filterMode = cudaFilterModeLinear;
    tex.normalized = true;
    cudaBindTextureToArray(tex, cuArray, channelDesc);

    // 预跑
    transformKernel<<<dimGrid, dimBlock, 0>>>(d_data, width, height, angle);
    cudaDeviceSynchronize();

    StopWatchInterface *timer = NULL;
    sdkCreateTimer(&timer);
    sdkStartTimer(&timer);
    
    transformKernel<<<dimGrid, dimBlock, 0>>>(d_data, width, height, angle);
    
    cudaDeviceSynchronize();
    sdkStopTimer(&timer); 
    sdkDeleteTimer(&timer);
    printf("
	Cost time: %f ms, %.2f Mpixels/sec
", sdkGetTimerValue(&timer), (width *height / (sdkGetTimerValue(&timer) / 1000.0f)) / 1e6);

    // 结果回收、输出和检验
    cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
    sdkSavePGM("D:\Code\CUDA\cudaProjectTemp\data\output.pgm", h_data, width, height);
    printf("
	Save output file.
");
    printf("
	Finish, return %s.
", compareData(h_data, h_dataRef, width * height, MIN_EPSILON_ERROR, 0.0f) ? "Passed" : "Failed");

    cudaFree(d_data);
    cudaFreeArray(cuArray);
    getchar();
    return 0;
}

▶ 输出结果

Start.

Device GeForce GTX 1070, Multi-Processors: 16, SM 6.1

Load input files, 512 x 512 pixels

Cost time: 0.000000 ms, inf Mpixels/sec

Save output file.

Finish, return Passed

▶ 涨姿势

● 使用函数 sdkLoadPGM() 读取图片数据

// helper_image.h
inline bool __loadPPM(const char *file, unsigned char **data, unsigned int *w, unsigned int *h, unsigned int *channels)
{
    FILE *fp = NULL;
    if (FOPEN_FAIL(FOPEN(fp, file, "rb")))
    {
        std::cerr << "__LoadPPM() : Failed to open file: " << file << std::endl;
        return false;
    }

    // check header
    char header[PGMHeaderSize];
    if (fgets(header, PGMHeaderSize, fp) == NULL)
    {
        std::cerr << "__LoadPPM() : reading PGM header returned NULL" << std::endl;
        return false;
    }
    if (strncmp(header, "P5", 2) == 0)
    {
        *channels = 1;
    }
    else if (strncmp(header, "P6", 2) == 0)
    {
        *channels = 3;
    }
    else
    {
        std::cerr << "__LoadPPM() : File is not a PPM or PGM image" << std::endl;
        *channels = 0;
        return false;
    }

    // parse header, read maxval, width and height
    unsigned int width = 0;
    unsigned int height = 0;
    unsigned int maxval = 0;
    unsigned int i = 0;
    while (i < 3)
    {
        if (fgets(header, PGMHeaderSize, fp) == NULL)
        {
            std::cerr << "__LoadPPM() : reading PGM header returned NULL" << std::endl;
            return false;
        }

        if (header[0] == '#')
        {
            continue;
        }

        if (i == 0)
        {
            i += SSCANF(header, "%u %u %u", &width, &height, &maxval);
        }
        else if (i == 1)
        {
            i += SSCANF(header, "%u %u", &height, &maxval);
        }
        else if (i == 2)
        {
            i += SSCANF(header, "%u", &maxval);
        }
    }

    // check if given handle for the data is initialized
    if (NULL != *data)
    {
        if (*w != width || *h != height)
        {
            std::cerr << "__LoadPPM() : Invalid image dimensions." << std::endl;
        }
    }
    else
    {
        *data = (unsigned char *)malloc(sizeof(unsigned char) * width * height **channels);
        *w = width;
        *h = height;
    }

    // read and close file
    if (fread(*data, sizeof(unsigned char), width * height **channels, fp) == 0)
    {
        std::cerr << "__LoadPPM() read data returned error." << std::endl;
    }

    fclose(fp);
    return true;
}

template <class T> inline bool sdkLoadPGM(const char *file, T **data, unsigned int *w, unsigned int *h)
{
    unsigned char *idata = NULL;
    unsigned int channels;

    if (!__loadPPM(file, &idata, w, h, &channels))
        return false;
    unsigned int size = *w **h * channels;

    if (*data == NULL)// 如果 T **data 没有初始化，则按照读取的 size 进行初始化 
        *data = (T *)malloc(sizeof(T) * size);

    std::transform(idata, idata + size, *data, ConverterFromUByte<T>());// 拷贝数据到 data 中

    free(idata);
    return true;
}

● 使用到的表面写入函数原型

// surface_functions.h
template<class T> static __device__ __forceinline__ void surf2Dwrite(T val, surface<void, cudaSurfaceType2D> surf, int x, int y, enum cudaSurfaceBoundaryMode mode = cudaBoundaryModeTrap)
{
#ifdef __CUDA_ARCH__ 
    __nv_tex_surf_handler("__surf2Dwrite_v2", (typename __nv_surf_trait<T>::cast_type)&val, (int)sizeof(T), surf, x, y, mode);
#endif
}