10.7 一维表面内存的读写

流处理器簇2.0出现之前，CUDA内核只能通过纹理操作对CUDA数组内容进行访问。至于其他CUDA数组访问方式，包括所有的写访问，都只能通过如CUDAMemcpyToArray()之类的内存复制函数进行。对于CUDA内核而言，对指定内存区域进行纹理读取以及写操作的唯一方式就是将纹理引用与线性设备内存进行绑定。

但若使用流处理器簇2.x上新增加的表面内存读写函数，开发者可以将CUDA数组与表面引用进行绑定，并调用surf1Dread()与surf1Dwrite()指令，对内核中的CUDA数组执行读写操作。与拥有特有硬件缓存的纹理读取不同，表面内存的读写操作将通过二级缓存实现，与进行全局加载和存储时类似。

提示 为了将表面引用绑定到CUDA数组，CUDA数组必须以CUDAArraySurfaceLoadStore的标志进行创建。

一维表面内存读写指令的声明如下所示:

template<class Type> Type surf1Dread(surface<void, l> surfRef, int x, boundaryMode =EMENTBoundaryModeTrap);
template<class Type> void surf1Dwrite(Type data, surface<void, l>
    surfRef, int x, boundaryMode =EMENTBoundaryModeTrap);

从上述代码可以看出，这些指令并不是强类型的（type-strong），表面引用是以void类型声明的，调用surf1Dread()或surf1Dwrite()时内存事务的大小取决于sizeof（Type）。偏移x的单位是字节，且必须按照sizeof（Type）对齐。对于如int或float类型的4字节操作数而言，该偏移必须能被4整除，而对于short类型而言，其必须能被2整除。其他类型以此类推。

表面读取支持的功能远远少于纹理操作的功能 [1]。其只支持非格式化的读写操作，而且没有类型转换或插值功能，边界处理的模式只有两种。

表面读写的边界处理方式与纹理读取的不同。对于纹理而言，该操作受纹理引用中的寻址模式控制。而对于表面读写而言，越界偏移值的处理方式由surf1Dread()或surf1Dwrite()中的一个参数指定。越界索引可导致两种情况：若将surf1Dread()的参数设置为CUDABoundModeTrap，处理越界索引时将抛出硬件异常；而将surf1Dwrite()的参数设置为CUDABoundaryModeZero时，surf1Dread()函数读得的值为0，surf1Dwrite()函数将被忽略。

由于表面引用的无类型的特点，很容易就能写出一个一维的针对所有类型的memset模板函数。

surface <void, l> surf1D;
template <typename T>
    global void
surf1Dmemset(int index, T value, size_t N)
{
    for (size_t i = blockIdx.x*blockDim.x + threadIdx.x; i < N; i += blockDim.x*gridDim.x)
        surf1Dwrite(value, surf1D, (index+i)*sizeof(T));
}

该内核存在于案例surf1Dmemset.cu中，为演示之用，该程序创建了一个64字节的CUDA数组，并使用上述内核初始化，最后以浮点和整型数的形式打印数组。

以下是一个通用的主机端函数模板代码，其将 CUDABindSurfaceToArray()与该内核的调用封装到一起。

template<typename T>   
cudaError_t   
surf1Dmemset( cudaArray *array, int offset, T value, size_t N)   
{   
    cudaError_t status;   
    CUDA_CHECK(cudaBindSurfaceToArray surf1D, array);   
    surf1Dmemset_kernel << 2,384 >> (0, value, 4*NUM_VALUE);   
}

纹理引用的无类型特性使得其模板结构的实现要比纹理的实现简单很多。由于纹理引用既是强类型又是全局的，因此其不能通过参数列表模板化为通用的函数。若将该函数的一行调用代码由

CUDART_CHECK(surf1Dmemset(array, 0, 3.141592654f, NUM_VALUES));

改为

CUDART_CHECK(surf1Dmemset(array, 0, (short) Oxbeef, 2*NUM_VALUES));

0x40490fdb 0x40490fdb ... (16 times)  
3.141593E+00 3.141593E+00 ... (16 times)

变为

Oxbeefbeef Oxbeefbeef ... (16 times) -4.68253E-01 -4.68253E-01 ... (16 times)

[1] 事实上, CUDA可以利用直接对CUDA数组操作的指令而彻底地绕过表面引用的实现。表面引用主要是为了与纹理引用相区别, 以提供与基于每个指令不同的基于每个表面引用的行为。