CUDA texture surface应用

最新推荐文章于 2024-10-08 20:08:41 发布

fb_help

最新推荐文章于 2024-10-08 20:08:41 发布

阅读量891

点赞数

分类专栏： CUDA 文章标签： cuda texture surface

本文链接：https://blog.csdn.net/fb_help/article/details/88085868

版权

CUDA 专栏收录该内容

51 篇文章

订阅专栏

本文探讨了CUDA中纹理和表面内存的使用及其对性能的影响。纹理提供了插值功能，适用于需要浮点坐标读取的应用场景，而表面内存则允许更灵活的内存访问模式。通过绑定纹理和表面到设备内存，可以利用硬件加速和缓存机制，提高数据访问速度，尤其对于重复访问模式，如卷积计算，能显著提升效率。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

CUDA texture surface应用

纹理的优势

使用纹理有几个性能优势。纹理可以暗示插值（即，使用浮点坐标从纹理读取）。任何需要这种数据插值的应用程序都可以受益于GPU上纹理单元内的HW插值引擎。
在任意GPU代码中使用纹理最重要的另一个好处是纹理缓存，它可以备份存储在全局内存中的纹理。纹理是一种只读操作，但如果您有一个只读数据数组，纹理缓存可能会改善或扩展您快速访问数据的能力。这通常意味着您的函数中必须存在数据局部性/数据重用，这些函数访问存储在纹理机制中的数据。检索到的纹理数据不会破坏L1缓存中的任何内容，因此通常这种类型的数据分段/优化将成为围绕数据缓存的更大策略的一部分。如果L1高速缓存上没有其他请求，则纹理机制/高速缓存不能提供比已经在L1中更快的数据访问。

即
纹理适合插值（纹理是以float坐标拾取的）
纹理对多次访问或局部访问的模式友好，例如说图像局部窗口（如卷积）计算

实例

#include <stdio.h>
#include <thrust\device_vector.h>

// --- 2D float texture
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

// --- 2D surface memory
surface<void, 2> surf2D;

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

/*************************************/
/* cudaArray PRINTOUT TEXTURE KERNEL */
/*************************************/
__global__ void cudaArrayPrintoutTexture(int width, int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    printf("Thread index: (%i, %i); cudaArray = %f\n", x, y, tex2D(texRef, x / (float)width + 0.5f, y / (float)height + 0.5f));
}

/*************************************/
/* cudaArray PRINTOUT TEXTURE KERNEL */
/*************************************/
__global__ void cudaArrayPrintoutSurface(int width, int height)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    float temp;

    surf2Dread(&temp, surf2D, x * 4, y);

    printf("Thread index: (%i, %i); cudaArray = %f\n", x, y, temp);
}

/********/
/* MAIN */
/********/
void main()
{
    int width = 3, height = 3;

    thrust::host_vector<float> h_data(width*height, 3.f);

    // --- Allocate CUDA array in device memory
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

    cudaArray* cuArray;

    /*******************/
    /* TEXTURE BINDING */
    /*******************/
    gpuErrchk(cudaMallocArray(&cuArray, &channelDesc, width, height));

    // --- Copy to host data to device memory
    gpuErrchk(cudaMemcpyToArray(cuArray, 0, 0, thrust::raw_pointer_cast(h_data.data()), width*height*sizeof(float), cudaMemcpyHostToDevice));

    // --- Set texture parameters
    texRef.addressMode[0] = cudaAddressModeWrap;
    texRef.addressMode[1] = cudaAddressModeWrap;
    texRef.filterMode = cudaFilterModeLinear;
    texRef.normalized = true;

    // --- Bind the array to the texture reference
    gpuErrchk(cudaBindTextureToArray(texRef, cuArray, channelDesc));

    // --- Invoking printout kernel
    dim3 dimBlock(3, 3);
    dim3 dimGrid(1, 1);
    cudaArrayPrintoutTexture<<<dimGrid, dimBlock>>>(width, height);

    gpuErrchk(cudaUnbindTexture(texRef));

    gpuErrchk(cudaFreeArray(cuArray));

    /******************/
    /* SURFACE MEMORY */
    /******************/
    gpuErrchk(cudaMallocArray(&cuArray, &channelDesc, width, height, cudaArraySurfaceLoadStore));

    // --- Copy to host data to device memory
    gpuErrchk(cudaMemcpyToArray(cuArray, 0, 0, thrust::raw_pointer_cast(h_data.data()), width*height*sizeof(float), cudaMemcpyHostToDevice));

    gpuErrchk(cudaBindSurfaceToArray(surf2D, cuArray));

    cudaArrayPrintoutSurface<<<dimGrid, dimBlock>>>(width, height);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaFreeArray(cuArray));
}