0_Simple__simplePitchLinearTexture

对比设备线性二维数组和 CUDA 二维数组在纹理引用中的效率

▶ 源代码。分别绑定相同大小的设备线性二维数组和 CUDA 二维数组为纹理引用，做简单的平移操作，重复若干次计算带宽和访问速度。

#include <stdio.h>
#ifdef _WIN32
#  define WINDOWS_LEAN_AND_MEAN
#  define NOMINMAX
#  include <windows.h>
#endif
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_functions.h>
#include <helper_cuda.h>

#define NUM_REPS 100  // test 重复次数
#define TILE_DIM 16   // 线程块尺寸

texture<float, 2, cudaReadModeElementType> texRefPL;
texture<float, 2, cudaReadModeElementType> texRefArray;

__global__ void shiftPitchLinear(float *odata, int pitch, int width, int height, int shiftX, int shiftY)
{
    int xid = blockIdx.x * blockDim.x + threadIdx.x;
    int yid = blockIdx.y * blockDim.y + threadIdx.y;

    odata[yid * pitch + xid] = tex2D(texRefPL, (xid + shiftX) / (float)width, (yid + shiftY) / (float)height);
}

__global__ void shiftArray(float *odata, int pitch, int width, int height, int shiftX, int shiftY)
{
    int xid = blockIdx.x * blockDim.x + threadIdx.x;
    int yid = blockIdx.y * blockDim.y + threadIdx.y;

    odata[yid * pitch + xid] = tex2D(texRefArray, (xid + shiftX) / (float)width, (yid + shiftY) / (float)height);
}

bool test()
{
    bool result = true;
    int i, j, ishift, jshift;
    // 数组大小以及 x，y 方向上的偏移量
    const int nx = 2048;
    const int ny = 2048;
    const int x_shift = 5;
    const int y_shift = 7;
    if ((nx % TILE_DIM) || (ny % TILE_DIM))
    {
        printf("nx and ny must be multiples of TILE_DIM
");
        return EXIT_FAILURE;
    }
    dim3 dimGrid(nx / TILE_DIM, ny / TILE_DIM), dimBlock(TILE_DIM, TILE_DIM);

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    //int devID = findCudaDevice(argc, (const char **)argv);// 使用device 0，不再使用命令行参数进行判断
    
    // 申请内存
    float *h_idata = (float *)malloc(sizeof(float) * nx * ny);
    float *h_odata = (float *)malloc(sizeof(float) * nx * ny);
    float *h_ref = (float *)malloc(sizeof(float) * nx * ny);
    for (int i = 0; i < nx * ny; ++i)
        h_idata[i] = (float)i;
    float *d_idataPL;
    size_t d_pitchBytes;
    cudaMallocPitch((void **)&d_idataPL, &d_pitchBytes, nx * sizeof(float), ny);
    cudaArray *d_idataArray;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    cudaMallocArray(&d_idataArray, &channelDesc, nx, ny);
    float *d_odata;
    cudaMallocPitch((void **)&d_odata, &d_pitchBytes, nx * sizeof(float), ny);

    // 拷贝内存（两组）
    size_t h_pitchBytes = nx * sizeof(float);
    cudaMemcpy2D(d_idataPL, d_pitchBytes, h_idata, h_pitchBytes, nx * sizeof(float), ny, cudaMemcpyHostToDevice);
    cudaMemcpyToArray(d_idataArray, 0, 0, h_idata, nx * ny * sizeof(float), cudaMemcpyHostToDevice);

    // 绑定纹理（两组）
    texRefPL.normalized = 1;
    texRefPL.filterMode = cudaFilterModePoint;
    texRefPL.addressMode[0] = cudaAddressModeWrap;
    texRefPL.addressMode[1] = cudaAddressModeWrap;
    cudaBindTexture2D(0, &texRefPL, d_idataPL, &channelDesc, nx, ny, d_pitchBytes);

    texRefArray.normalized = 1;
    texRefArray.filterMode = cudaFilterModePoint;
    texRefArray.addressMode[0] = cudaAddressModeWrap;
    texRefArray.addressMode[1] = cudaAddressModeWrap;
    cudaBindTextureToArray(texRefArray, d_idataArray, channelDesc);

    // 理论计算结果
    for (i = 0; i < ny; i++)
    {
        for (j = 0; j < nx; ++j)
            h_ref[i * nx + j] = h_idata[(i + y_shift) % ny * nx + (j + x_shift) % nx];
    }

    // 使用线性数组的纹理计算
    cudaMemset2D(d_odata, d_pitchBytes, 0, nx * sizeof(float), ny);
    cudaEventRecord(start, 0);
    for (int i = 0; i < NUM_REPS; ++i)
        shiftPitchLinear << <dimGrid, dimBlock >> > (d_odata, (int)(d_pitchBytes / sizeof(float)), nx, ny, x_shift, y_shift);
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float timePL;
    cudaEventElapsedTime(&timePL, start, stop);

    // 检查结果
    cudaMemcpy2D(h_odata, h_pitchBytes, d_odata, d_pitchBytes, nx * sizeof(float), ny, cudaMemcpyDeviceToHost);
    if (!compareData(h_ref, h_odata, nx*ny, 0.0f, 0.15f))
    {
        printf("
	 ShiftPitchLinear failed
");
        result = false;
    }

    // 使用 CUDA数组的纹理计算
    cudaMemset2D(d_odata, d_pitchBytes, 0, nx * sizeof(float), ny);
    cudaEventRecord(start, 0);
    for (int i = 0; i < NUM_REPS; ++i)
        shiftArray << <dimGrid, dimBlock >> > (d_odata, (int)(d_pitchBytes / sizeof(float)), nx, ny, x_shift, y_shift);
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float timeArray;
    cudaEventElapsedTime(&timeArray, start, stop);

    // 检查结果
    cudaMemcpy2D(h_odata, h_pitchBytes, d_odata, d_pitchBytes, nx * sizeof(float), ny, cudaMemcpyDeviceToHost);
    if (!compareData(h_ref, h_odata, nx*ny, 0.0f, 0.15f))
    {
        printf("
	ShiftArray failed
");
        result = false;
    }

    // 计算带宽和读取速度
    float bandwidthPL = 2.f * nx * ny * sizeof(float) / (timePL / 1000.f / NUM_REPS * 1.e+9f);
    float bandwidthArray = 2.f * nx * ny * sizeof(float) / (timeArray / 1000.f / NUM_REPS * 1.e+9f);
    printf("
	Bandwidth for pitch linear: %.2f GB/s; for array: %.2f GB/s
", bandwidthPL, bandwidthArray);
    float fetchRatePL = nx * ny / 1.e+6f / (timePL / 1000.0f / NUM_REPS); 
    float fetchRateArray = nx * ny / 1.e+6f / (timeArray / 1000.0f / NUM_REPS); 
    printf("
	Texture fetch rate for pitch linear: %.2f Mpix/s; for array: %.2f Mpix/s
", fetchRatePL, fetchRateArray);

    // 回收工作
    free(h_idata);
    free(h_odata);
    free(h_ref);
    cudaUnbindTexture(texRefPL);
    cudaUnbindTexture(texRefArray);
    cudaFree(d_idataPL);
    cudaFreeArray(d_idataArray);
    cudaFree(d_odata);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    return result;
}

int main(int argc, char **argv)
{
    printf("
	Start
");
    printf("
	Finished, %s
", test() ? "Passed" : "Failed");

    getchar();
    return 0;
}

▶ 输出结果

    Start

    Bandwidth for pitch linear: 12.58 GB/s; for array: 14.64 GB/s

    Texture fetch rate for pitch linear: 1573.09 Mpix/s; for array: 1829.39 Mpix/s

    Finished, Passed

▶ 涨姿势

● 用到的函数都在以前的，有关线性二维数组和纹理内存使用方法的博客汇总讨论过了。

● 由运行结果可知，使用二维纹理引用时，CUDA 二维数组的效率比线性二维数组更高。