LBM base class design

[kks32]

How should we structure the LBM base class with distribution functions f defined on GPU?

One possibility is to have the distribution function f be a flattened array that gets 3 dimensions (lx, ly, and lz) in the constructor. As described here: https://stackoverflow.com/questions/12924155/sending-3d-array-to-cuda-kernel/12925014#12925014

This paper also uses flattened array: index new = 𝑥id + 𝑦id × 𝑁𝑥 + 𝑧id × 𝑁𝑥 × 𝑁𝑦

The alternative options are:

• Pitched array https://stackoverflow.com/questions/6137218/how-can-i-add-up-two-2d-pitched-arrays-using-nested-for-loops
• 3D/4D pointers

@Amber1995 @reihos What are the design advantages/disadvantages of one over the other?

[Amber1995]

There are some pros and cons I could think up with using different structures of distribution arrays:

1. Flattened array:

• Pros：contiguous memory, request for relatively small memory size compared to 3d/4d pointers, the most conventional approach and thus enough references
• Cons：relativley complicated semantics, not as straightforward as 3d/4d arrays

2. Pitched array
   I don't have much knowledge to this but I read through the CUDA documentation and others' examples and concluded several points as below

• Pros: It is supposed to be a standardized built-in functions and thus more convinient to create. For example, just use "cudaPitchedPtr devPitchedPtr;
  cudaMalloc3D(&devPitchedPtr, extent);"
• Cons: Little knowledge on establishing 4D array since there is no function named cudaMalloc4D;
  It's obscure how to iterate through the array or locate an specific element at (x,y,z) using slice and pitch.

3. 3D/4D array(only dynamic allocated array is discussed here)

• Pros: no necessary transition from the treadID, blockID and blockDimension to the coordination of the element; perfectly match the geometry of GPU architecture;
• Cons: If it's dynamically allocated, memories are requested to store the address of lower rank pointers apart from the lowest rank pointer. For example, for a 3d array, the top rank mallocates lxsizeof(real**), the second rank mallocates lxlysizeof(real), and finally the lowest rank mallocates lxlylz*sizeof(real); however, the number of bytes requested by the first two ranks is not wanted.
  Another disadvantage is that when copying between the host and device, it is not allowed to copy the top ranked pointer because the memory is not linearly arranged. You should do something like a nested for loop to copy the bottom ranked pointer.

[Amber1995]

So, I would definitely support the flattened array.

[kks32]

@Amber1995 do you wanna create a RFC for the LBM base class with the flattened array then?

[Amber1995]

Of course. I was expecting to try flattened array for this time.