CUDA学习（二十）

表面记忆：
对于计算能力为2.0或更高的设备，使用cudaArraySurfaceLoadStore标志创建的CUDA数组（可在Cubemap表面中描述）可以使用Surface函数中描述的函数通过曲面对象或曲面参考进行读写。
表14列出了根据设备计算能力的最大表面宽度，高度和深度。

表面对象API：
使用struct cudaResourceDesc类型的资源描述中的cudaCreateSurfaceObject（）创建表面对象。
以下代码示例将一些简单的转换内核应用于纹理：

// Simple copy kernel
__global__ void copyKernel(cudaSurfaceObject_t inputSurfObj,
    cudaSurfaceObject_t outputSurfObj,
    int width, int height)
{
    // Calculate surface coordinates
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        uchar4 data;
        // Read from input surface
        surf2Dread(&data, inputSurfObj, x * 4, y);
        // Write to output surface
        surf2Dwrite(data, outputSurfObj, x * 4, y);
    }
}
// Host code
int main()
{
    // Allocate CUDA arrays in device memory
    cudaChannelFormatDesc channelDesc =
        cudaCreateChannelDesc(8, 8, 8, 8,
            cudaChannelFormatKindUnsigned);
    cudaArray* cuInputArray;
    cudaMallocArray(&cuInputArray, &channelDesc, width, height,
        cudaArraySurfaceLoadStore);
    cudaArray* cuOutputArray;
    cudaMallocArray(&cuOutputArray, &channelDesc, width, height,
        cudaArraySurfaceLoadStore);
    // Copy to device memory some data located at address h_data
    // in host memory
    cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size,
        cudaMemcpyHostToDevice);
    // Specify surface
    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;
    // Create the surface objects
    resDesc.res.array.array = cuInputArray;
    cudaSurfaceObject_t inputSurfObj = 0;
    cudaCreateSurfaceObject(&inputSurfObj, &resDesc);
    resDesc.res.array.array = cuOutputArray;
    cudaSurfaceObject_t outputSurfObj = 0;
    cudaCreateSurfaceObject(&outputSurfObj, &resDesc);
    // Invoke kernel
    dim3 dimBlock(16, 16);
    dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x,
        (height + dimBlock.y - 1) / dimBlock.y);
    copyKernel << <dimGrid, dimBlock >> >(inputSurfObj,
        outputSurfObj,
        width, height);
    // Destroy surface objects
    cudaDestroySurfaceObject(inputSurfObj);
    cudaDestroySurfaceObject(outputSurfObj);
    // Free device memory
    cudaFreeArray(cuInputArray);
    cudaFreeArray(cuOutputArray);
    return 0;
}

表面参考API：
表面引用在文件范围内声明为表面类型的变量：

surface<void, Type> surfRef;

其中Type指定曲面引用的类型，并且等于cudaSurfaceType1D，cudaSurfaceType2D，cudaSurfaceType3D，cudaSurfaceTypeCubemap，cudaSurfaceType1DLayered，cudaSurfaceType2DLayered或cudaSurfaceTypeCubemapLayered; Type是一个可选参数，默认为cudaSurfaceType1D。 表面引用只能被声明为静态全局变量，不能作为参数传递给函数。
在内核可以使用表面引用来访问CUDA数组之前，必须使用cudaBindSurfaceToArray（）将表面引用绑定到CUDA数组。
以下代码示例将表面引用绑定到CUDA阵列cuArray：
使用低级API：

surface<void, cudaSurfaceType2D> surfRef;
surfaceReference* surfRefPtr;
cudaGetSurfaceReference(&surfRefPtr, "surfRef");
cudaChannelFormatDesc channelDesc;
cudaGetChannelDesc(&channelDesc, cuArray);
cudaBindSurfaceToArray(surfRef, cuArray, &channelDesc);

使用高级API：

surface<void, cudaSurfaceType2D> surfRef;
cudaBindSurfaceToArray(surfRef, cuArray);

CUDA数组必须使用匹配维度和类型的表面函数以及匹配维数的表面引用来读写。 否则，读取和写入CUDA数组的结果是不确定的。
与纹理内存不同，表面内存使用字节寻址。 这意味着用于通过纹理函数访问纹理元素的x坐标需要乘以元素的字节大小才能通过曲面函数访问相同的元素。 例如，通过texRef使用tex1d（texRef，x）读取绑定到纹理参考texRef和表面参考surfRef的一维浮点CUDA数组的纹理坐标x上的元素，但是surf1Dread（surfRef，4 x ）通过surfRef。 类似地，通过texRef使用tex2d（texRef，x，y）来访问绑定到纹理参考texRef和表面参考surfRef的二维浮点CUDA数组的纹理坐标x和y处的元素，但是surf2Dread（surfRef，4 x，y）通过surfRef（y坐标的字节偏移是从CUDA阵列的底层行间距内部计算的）。
以下代码示例将一些简单的转换内核应用于纹理：

// 2D surfaces
surface<void, 2> inputSurfRef;
surface<void, 2> outputSurfRef;
// Simple copy kernel
__global__ void copyKernel(int width, int height)
{
// Calculate surface coordinates
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
uchar4 data;
// Read from input surface
surf2Dread(&data, inputSurfRef, x * 4, y);
// Write to output surface
surf2Dwrite(data, outputSurfRef, x * 4, y);
}
}
// Host code
int main()
{
// Allocate CUDA arrays in device memory
cudaChannelFormatDesc channelDesc =
cudaCreateChannelDesc(8, 8, 8, 8,
cudaChannelFormatKindUnsigned);
cudaArray* cuInputArray;
cudaMallocArray(&cuInputArray, &channelDesc, width, height,
cudaArraySurfaceLoadStore);
cudaArray* cuOutputArray;
cudaMallocArray(&cuOutputArray, &channelDesc, width, height,
cudaArraySurfaceLoadStore);
// Copy to device memory some data located at address h_data
// in host memory
cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size,
cudaMemcpyHostToDevice);
// Bind the arrays to the surface references
cudaBindSurfaceToArray(inputSurfRef, cuInputArray);
cudaBindSurfaceToArray(outputSurfRef, cuOutputArray);
// Invoke kernel
dim3 dimBlock(16, 16);
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x,
(height + dimBlock.y - 1) / dimBlock.y);
copyKernel<<<dimGrid, dimBlock>>>(width, height);
// Free device memory
cudaFreeArray(cuInputArray);
cudaFreeArray(cuOutputArray);
return 0;
}