Project

General

Profile

Task #280541

multi-dimensional CUDA arrays

Added by Felix Höfling about 3 years ago.

Status:
New
Priority:
Normal
Assignee:
-
Target version:
-
Start date:
2014-03-14
Due date:
% Done:

0%


Description

Many kernels receive a data pointer and separately the array size as int. 2D-Arrays in addition have the number of columns or stride/pitch. It would be merely cosmetics, but nice: combine these parameters in a small struct, something like a minimal CUDA-enable version of boost::multi_array (but with contiguous storage). Thereby, the call signatures of many kernels could be simplified (see neighbours/from_binning_kernel.cu for an example).

namespace cuda {

template <typename T, int rank>
struct multi_array {
    T* data;
    size_t size[rank];
    size_t stride[rank-1];
};

}

For rank=1, such a class exists internally: cuda::vector::container (in cuda_wrapper). Cell and neighbour lists are of rank 2, the array holding the stress tensor per particle would be an example for a rank-3 case.

In this context, we may evaluate the function cudaMallocPitch():
http://stackoverflow.com/questions/1047369/allocate-2d-array-on-device-memory-in-cuda

but the performance impact might be marginal:
http://devblogs.nvidia.com/parallelforall/how-access-global-memory-efficiently-cuda-c-kernels

Also available in: Atom PDF