For devices of compute capability 2.0 and higher, a CUDA array, created with the cudaArraySurfaceLoadStore flag, can be read and written via a surface object or surface reference using the functions described in Surface Functions.
On the other hand, a CUDA array, allocated with the flag cudaArrayLayered can have a layered structure.
Is it also possible to have CUDA array with both flags, so that it can be written to in a specific kernel and accessed (layeredwise) in a second kernel?
Yes. In all cases you would need to use the layered access functions, eg., you would need to use surf?Dlayeredwrite() to write to it, and similarly to read from it.
minimal example, 2D layered:
$ cat t2106.cu
#include <helper_cuda.h>
const int layers = 2;
const int xdim = 3;
const int ydim = 3;
__device__ float d_arr[xdim*ydim*layers] = {1.0, 1.1, 1.2, 1.3, 1.4, 1.5, 1.6, 1.7, 1.8, 2.0, 2.1, 2.2, 2.3, 2.4, 2.5, 2.6, 2.7, 2.8};
__global__ void write_kernel(cudaSurfaceObject_t surf)
{
for(int i = 0; i < layers; i++){
for(int j = 0; j < xdim; j++){
for(int k = 0; k < ydim; k++){
surf2DLayeredwrite<float>(d_arr[i*xdim*ydim+j*ydim+k], surf, j*sizeof(float), k, i);
}
}
}
}
__global__ void read_kernel(cudaSurfaceObject_t surf)
{
for(int i = 0; i < layers; i++){
for(int j = 0; j < xdim; j++){
for(int k = 0; k < ydim; k++){
printf(" layer:%d x:%d y:%d val: %f\n", i, j, k, surf2DLayeredread<float>(surf, j*sizeof(float), k, i));
}
}
}
}
int main(int argc, char **argv)
{
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaArray *dev_cu_array;
cudaSurfaceObject_t surf;
cudaResourceDesc res_desc;
memset(&res_desc, 0, sizeof(res_desc));
cudaExtent extent = make_cudaExtent(xdim, ydim, layers);
checkCudaErrors(cudaMalloc3DArray(&dev_cu_array,
&channelDesc,
extent,
cudaArrayLayered|cudaArraySurfaceLoadStore));
res_desc.resType = cudaResourceTypeArray;
res_desc.res.array.array = dev_cu_array;
checkCudaErrors(cudaCreateSurfaceObject(&surf, &res_desc));
write_kernel<<<1,1>>>(surf);
read_kernel<<<1,1>>>(surf);
checkCudaErrors(cudaDeviceSynchronize());
checkCudaErrors(cudaDestroySurfaceObject(surf));
checkCudaErrors(cudaFreeArray(dev_cu_array));
}
$ nvcc -o t2106 t2106.cu -I/usr/local/cuda/samples/common/inc
$ compute-sanitizer ./t2106
========= COMPUTE-SANITIZER
layer:0 x:0 y:0 val: 1.000000
layer:0 x:0 y:1 val: 1.100000
layer:0 x:0 y:2 val: 1.200000
layer:0 x:1 y:0 val: 1.300000
layer:0 x:1 y:1 val: 1.400000
layer:0 x:1 y:2 val: 1.500000
layer:0 x:2 y:0 val: 1.600000
layer:0 x:2 y:1 val: 1.700000
layer:0 x:2 y:2 val: 1.800000
layer:1 x:0 y:0 val: 2.000000
layer:1 x:0 y:1 val: 2.100000
layer:1 x:0 y:2 val: 2.200000
layer:1 x:1 y:0 val: 2.300000
layer:1 x:1 y:1 val: 2.400000
layer:1 x:1 y:2 val: 2.500000
layer:1 x:2 y:0 val: 2.600000
layer:1 x:2 y:1 val: 2.700000
layer:1 x:2 y:2 val: 2.800000
========= ERROR SUMMARY: 0 errors
$