使用READ和WRITE构造纹理内存

Question

我正在开发一个CUDA应用程序,其中内核必须多次访问全局内存.所有CTA随机访问此内存(没有位置,因此不能使用共享内存).我需要优化它.我听说纹理内存可以缓解这个问题,但内核可以读写纹理内存吗？1D纹理记忆？2D纹理记忆？还有CUDA阵列呢？

Answer 1

CUDA纹理是只读的.纹理读取被缓存.因此,性能提升是概率性的.

CUDA Toolkit 3.1以后也有可写纹理称为Surfaces,但它们仅适用于Compute Capability> = 2.0的设备.曲面就像纹理,但优点是它们也可以由内核编写.

曲面只能绑定到cudaArray使用flag创建cudaArraySurfaceLoadStore.

Answer 2

这是 sgarizvi 回答的后续。

如今，具有计算能力的卡>=2.0比在2012问这个问题时更常见。

下面是一个关于如何使用CUDA 表面内存写入纹理的最小示例。

#include <stdio.h>

#include "TimingGPU.cuh"
#include "Utilities.cuh"

surface<void, cudaSurfaceType1D> surfD;

/*******************/
/* KERNEL FUNCTION */
/*******************/
__global__ void SurfaceMemoryWrite(const int N) {

    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    surf1Dwrite((float)tid, surfD, tid * sizeof(float), cudaBoundaryModeTrap);
}

/********/
/* MAIN */
/********/
int main() {

    const int N = 10;

    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    //Alternatively
    //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

    cudaArray *d_arr;   gpuErrchk(cudaMallocArray(&d_arr, &channelDesc, N, 1, cudaArraySurfaceLoadStore));
    gpuErrchk(cudaBindSurfaceToArray(surfD, d_arr));

    SurfaceMemoryWrite<<<1, N>>>(N);

    float *h_arr = new float[N];
    gpuErrchk(cudaMemcpyFromArray(h_arr, d_arr, 0, 0, N * sizeof(float), cudaMemcpyDeviceToHost));

    for (int i=0; i<N; i++) printf("h_arr[%i] = %f\n", i, h_arr[i]);

    return 0;
}