Update the memory selection layout when there is a layout mismatch

[anagainaru]

Since we set the memory space after we set all the variable shapes and selections, the memory selection needs to be updated together with the variable shape.

        const adios2::Dims shape{static_cast<size_t>(Nx * mpiSize), Ny};
        const adios2::Dims start{static_cast<size_t>(Nx * mpiRank), 0};
        const adios2::Dims count{Nx, Ny};
        auto var = io.DefineVariable<float>("r32", shape, start, count);

        const adios2::Dims memoryStart = {ghostCells, ghostCells};
        const adios2::Dims memoryCount = {totalNx, totalNy};
        var.SetMemorySelection({memoryStart, memoryCount});

        var.SetMemorySpace(adios2::MemorySpace::GPU);

The SetMemorySpace call was updating the shape, start and count but not the memory selection.

@franzpoeschel this is why your code was writing the wrong data for your 2D array. Could you please double check?

[franzpoeschel]

Hello Ana, thank you for your help.

Unfortunately, the reproducer linked in this comment still produces this file for me, even with this PR:

> bpls MemSelection.bp/ -d                                                                                                                                  
  double   cpu_pointer_with_mem_selection     {12, 12}                                                                                                      
    ( 0, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 0, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 1, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 1, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 2, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 2, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 3, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 3, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 4, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 4, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 5, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 5, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 6, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 6, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 7, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 7, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 8, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 8, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 9, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 9, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    (10, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    (10, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    (11, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    (11, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
                                                                                                                                                            
  double   cpu_pointer_without_mem_selection  {12, 12}                                                                                                      
    ( 0, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 0, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170                                                                
    ( 1, 0)    0 0 0 0 0 0                                                                                                                                  
    ( 1, 6)    0 0 0 0 0 0
    ( 2, 0)    0 0 0 0 0 0
    ( 2, 6)    0 0 0 0 0 0
    ( 3, 0)    0 0 0 0 0 0
    ( 3, 6)    0 0 0 0 0 0
    ( 4, 0)    0 0 0 0 0 0
    ( 4, 6)    0 0 0 0 0 0
    ( 5, 0)    0 0 0 0 5.85364e+170 5.85364e+170
    ( 5, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    ( 6, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 0 0
    ( 6, 6)    0 0 0 0 0 0
    ( 7, 0)    0 0 0 0 0 0
    ( 7, 6)    0 0 0 0 0 0
    ( 8, 0)    0 0 0 0 0 0
    ( 8, 6)    0 0 0 0 0 0
    ( 9, 0)    0 0 0 0 0 0
    ( 9, 6)    0 0 0 0 0 0
    (10, 0)    0 0 0 0 0 0
    (10, 6)    0 0 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    (11, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    (11, 6)    5.85364e+170 5.85364e+170 0 0 0 0

  double   gpu_pointer_with_mem_selection     {12, 12}
    ( 0, 0)    0 0 0 0 0 0
    ( 0, 6)    0 0 0 0 0 0
    ( 1, 0)    0 0 0 0 0 0
    ( 1, 6)    0 0 0 0 0 0
    ( 2, 0)    0 0 0 0 0 0
    ( 2, 6)    0 0 0 0 0 0
    ( 3, 0)    0 0 0 0 0 0
    ( 3, 6)    0 0 0 0 0 0
    ( 4, 0)    0 0 0 0 0 0
    ( 4, 6)    0 0 0 0 0 0
    ( 5, 0)    0 0 0 0 0 0
    ( 5, 6)    0 0 0 0 0 0
    ( 6, 0)    0 0 0 0 0 0
    ( 6, 6)    0 0 0 0 0 0
    ( 7, 0)    0 0 0 0 0 0
    ( 7, 6)    0 0 0 0 0 0
    ( 8, 0)    0 0 0 0 0 0
    ( 8, 6)    0 0 0 0 0 0
    ( 9, 0)    0 0 0 0 0 0
    ( 9, 6)    0 0 0 0 0 0
    (10, 0)    0 0 0 0 0 0
    (10, 6)    0 0 0 0 0 0
    (11, 0)    0 0 0 0 0 0
    (11, 6)    0 0 0 0 0 0

  double   gpu_pointer_without_mem_selection  {12, 12}
    ( 0, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    ( 0, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    ( 1, 0)    0 0 0 0 0 0
    ( 1, 6)    0 0 0 0 0 0
    ( 2, 0)    0 0 0 0 0 0
    ( 2, 6)    0 0 0 0 0 0
    ( 3, 0)    0 0 0 0 0 0
    ( 3, 6)    0 0 0 0 0 0
    ( 4, 0)    0 0 0 0 0 0
    ( 4, 6)    0 0 0 0 0 0
    ( 5, 0)    0 0 0 0 5.85364e+170 5.85364e+170
    ( 5, 6)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    ( 6, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 0 0
    ( 6, 6)    0 0 0 0 0 0
    ( 7, 0)    0 0 0 0 0 0
    ( 7, 6)    0 0 0 0 0 0
    ( 8, 0)    0 0 0 0 0 0
    ( 8, 6)    0 0 0 0 0 0
    ( 9, 0)    0 0 0 0 0 0
    ( 9, 6)    0 0 0 0 0 0
    (10, 0)    0 0 0 0 0 0
    (10, 6)    0 0 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    (11, 0)    5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170 5.85364e+170
    (11, 6)    5.85364e+170 5.85364e+170 0 0 0 0

I will try this on other systems next week, since you report that you cannot reproduce this.

[anagainaru]

Ok, so two things I notice when I look at your code:

1. You are using the same memory selection for both CPU and GPU but the layout between the two is different so the data is not accessed the same for both. If you change the selection for the GPU to:
   var.SetMemorySelection({{0, 0}, {pitch / sizeof(value_type), square_size}}); do you still see zeros?

2. You are using cudaMallocPitch which maybe doesn't allocate continuous memory, which may explain why I cannot reproduce what you are seeing. This I am not sure about, I will check and come back to you.

I really think your issue now is a layout mismatch. CudaMemCpy doesn't reorganize data so you'll have the same data but we access it differently depending on if it's on the CPU and GPU. I can talk later today on teams if you want

[franzpoeschel]

2. You are using cudaMallocPitch which maybe doesn't allocate continuous memory, which may explain why I cannot reproduce what you are seeing. This I am not sure about, I will check and come back to you.

Correct, this is the reason why I want to use a memory selection here in the first place.
cudaMallocPitch allocates a 2D array and pads each line up to some platform-dependent power of 2, and I want to be rid of the padding in the output.

I output four variables in order to verify that this runs as I expect:

1. gpu_pointer_without_mem_selection outputs the memory pointer and ignores that there is a padding. The output shows this: The actual 12 data items per line are padded up by zeroes to a physical line length of 64.
2. In order to verify that I'm applying the memory selection correctly, I download the data to CPU with a simple cudaMemcpy(), ignoring the memory layout, meaning that also the padding is downloaded.
   I verify this by writing the variable cpu_pointer_without_mem_selection: It has the exact same content as gpu_pointer_without_mem_selection, so the memory layouts on CPU and GPU are identical.
3. I then apply the memory selection on CPU and write this to cpu_pointer_with_mem_selection. This gives me the expected results.
4. Trying the same on the GPU in gpu_pointer_with_mem_selection gives me only zeroes.

1. You are using the same memory selection for both CPU and GPU but the layout between the two is different so the data is not accessed the same for both. If you change the selection for the GPU to:
   var.SetMemorySelection({{0, 0}, {pitch / sizeof(value_type), square_size}}); do you still see zeros?

I'll try that next week. Since you mention that you access CPU/GPU data differently, even when it has the same layout as in this case, this seems like I need to know some special tricks to output padded GPU data with memory selections?
Thank you for helping me to get this running!