Memory Selection for GPU pointers #4337

franzpoeschel · 2024-09-09T14:43:54Z

franzpoeschel
Sep 9, 2024

Hello,

I noticed that SetMemorySelection has no effect when dealing with GPU pointers, it rather leads to a dataset filled with all zeroes.
I have a 2-dimensional dataset on the GPU with extent [66, 66], padded to [66, 128]. I.e., when writing directly from the GPU buffer without specifying the memory selection, the data shows the padded regions:

> bpls heat_000001.bp5/ -d                                                                                             
  double   /data/1/images/heat  {66, 66}                                                                               
    ( 0, 0)    0 0.0483134 0.0965139 0.144489 0.192127 0.239316                                                        
    ( 0, 6)    0.285946 0.331908 0.377095 0.421401 0.464723 0.50696                                                    
    ( 0,12)    0.548013 0.587785 0.626185 0.663123 0.698511 0.732269                                                   
    ( 0,18)    0.764316 0.794578 0.822984 0.849468 0.873968 0.896427                                                   
    ( 0,24)    0.916792 0.935016 0.951057 0.964876 0.976441 0.985726
    ( 0,30)    0.992709 0.997373 0.999708 0.999708 0.997373 0.992709
    ( 0,36)    0.985726 0.976441 0.964876 0.951057 0.935016 0.916792
    ( 0,42)    0.896427 0.873968 0.849468 0.822984 0.794578 0.764316
    ( 0,48)    0.732269 0.698511 0.663123 0.626185 0.587785 0.548013
    ( 0,54)    0.50696 0.464723 0.421401 0.377095 0.331908 0.285946
    ( 0,60)    0.239316 0.192127 0.144489 0.0965139 0.0483134 1.22465e-16
    ( 1, 0)    0 0 0 0 0 0
    ( 1, 6)    0 0 0 0 0 0
    ( 1,12)    0 0 0 0 0 0
    ( 1,18)    0 0 0 0 0 0
    ( 1,24)    0 0 0 0 0 0
    ( 1,30)    0 0 0 0 0 0
    ( 1,36)    0 0 0 0 0 0
    ( 1,42)    0 0 0 0 0 0
    ( 1,48)    0 0 0 0 0 0
    ( 1,54)    0 0 0 0 0 0
    ( 1,60)    0 0 0.0483134 0.0966268 0.144827 0.192802
    ( 2, 0)    0.24044 0.287629 0.334259 0.380221 0.425408 0.469714
    ( 2, 6)    0.513037 0.555273 0.596326 0.636099 0.674499 0.711436
    ( 2,12)    0.746825 0.780582 0.812629 0.842891 0.871297 0.897781
    ( 2,18)    0.922281 0.94474 0.965106 0.98333 0.99937 1.01319
    ( 2,24)    1.02475 1.03404 1.04102 1.04569 1.04802 1.04802
    ( 2,30)    1.04569 1.04102 1.03404 1.02475 1.01319 0.99937
    ( 2,36)    0.98333 0.965106 0.94474 0.922281 0.897781 0.871297
    ( 2,42)    0.842891 0.812629 0.780582 0.746825 0.711436 0.674499
    ( 2,48)    0.636099 0.596326 0.555273 0.513037 0.469714 0.425408
    ( 2,54)    0.380221 0.334259 0.287629 0.24044 0.192802 0.144827
    ( 2,60)    0.0966268 0.0483134 0 0 0 0
    ( 3, 0)    0 0 0 0 0 0
    ( 3, 6)    0 0 0 0 0 0
    ( 3,12)    0 0 0 0 0 0
    ...

When I now specify the memory selection as [0, 0] -- [66, 128], the dataset shows only zeroes:

> bpls heat_000001.bp5/ -D
  double   /data/1/images/heat  {66, 66}
        step 0: 
          block 0: [ 0:65,  0:65]

> bpls heat_000001.bp5/ -d                                                                                             
  double   /data/1/images/heat  {66, 66}                                                                               
    ( 0, 0)    0 0 0 0 0 0                                                                                             
    ( 0, 6)    0 0 0 0 0 0                                                                                             
    ( 0,12)    0 0 0 0 0 0                                                                                             
    ( 0,18)    0 0 0 0 0 0                                                                                             
    ( 0,24)    0 0 0 0 0 0                                                                                             
    ( 0,30)    0 0 0 0 0 0                                                                                             
    ( 0,36)    0 0 0 0 0 0                                                                                             
    ( 0,42)    0 0 0 0 0 0                                                                                             
    ( 0,48)    0 0 0 0 0 0                                                                                             
    ( 0,54)    0 0 0 0 0 0                                                                                             
    ( 0,60)    0 0 0 0 0 0                                                                                             
    ( 1, 0)    0 0 0 0 0 0     
    …

The memory selection seems to be correct; when I download the dataset with padding to the CPU and store from there, the results look fine:

  double   /data/1/images/heat  {66, 66}
    ( 0, 0)    0 0.0483134 0.0965139 0.144489 0.192127 0.239316
    ( 0, 6)    0.285946 0.331908 0.377095 0.421401 0.464723 0.50696
    ( 0,12)    0.548013 0.587785 0.626185 0.663123 0.698511 0.732269
    ( 0,18)    0.764316 0.794578 0.822984 0.849468 0.873968 0.896427
    ( 0,24)    0.916792 0.935016 0.951057 0.964876 0.976441 0.985726
    ( 0,30)    0.992709 0.997373 0.999708 0.999708 0.997373 0.992709
    ( 0,36)    0.985726 0.976441 0.964876 0.951057 0.935016 0.916792
    ( 0,42)    0.896427 0.873968 0.849468 0.822984 0.794578 0.764316
    ( 0,48)    0.732269 0.698511 0.663123 0.626185 0.587785 0.548013
    ( 0,54)    0.50696 0.464723 0.421401 0.377095 0.331908 0.285946
    ( 0,60)    0.239316 0.192127 0.144489 0.0965139 0.0483134 1.22465e-16
    ( 1, 0)    0.0483134 0.0966268 0.144827 0.192802 0.24044 0.287629
    ( 1, 6)    0.334259 0.380221 0.425408 0.469714 0.513037 0.555273
    ( 1,12)    0.596326 0.636099 0.674499 0.711436 0.746825 0.780582
    ( 1,18)    0.812629 0.842891 0.871297 0.897781 0.922281 0.94474
    ( 1,24)    0.965106 0.98333 0.99937 1.01319 1.02475 1.03404
    ( 1,30)    1.04102 1.04569 1.04802 1.04802 1.04569 1.04102
    ( 1,36)    1.03404 1.02475 1.01319 0.99937 0.98333 0.965106
    ( 1,42)    0.94474 0.922281 0.897781 0.871297 0.842891 0.812629
    ( 1,48)    0.780582 0.746825 0.711436 0.674499 0.636099 0.596326
    ( 1,54)    0.555273 0.513037 0.469714 0.425408 0.380221 0.334259
    ( 1,60)    0.287629 0.24044 0.192802 0.144827 0.0966268 0.0483134
    ( 2, 0)    0.0965139 0.144827 0.193028 0.241003 0.288641 0.33583
    ( 2, 6)    0.38246 0.428422 0.473609 0.517915 0.561237 0.603474
    ( 2,12)    0.644526 0.684299 0.722699 0.759637 0.795025 0.828783
    ( 2,18)    0.86083 0.891092 0.919498 0.945982 0.970482 0.992941
    ( 2,24)    1.01331 1.03153 1.04757 1.06139 1.07295 1.08224
    ...

Is that a known restriction?

Answered by anagainaru

Sep 9, 2024

GPU pointers should work with memory selection. I'll try to reproduce this sometime this week and try to see what's going on, can you share a minimal code that I can use?

View full answer

anagainaru · 2024-09-09T15:11:38Z

anagainaru
Sep 9, 2024
Maintainer

GPU pointers should work with memory selection. I'll try to reproduce this sometime this week and try to see what's going on, can you share a minimal code that I can use?

17 replies

anagainaru Oct 7, 2024
Maintainer

I was finally able to reproduce this by using the CUDA backend instead of Kokkos on Perlmutter. Will return with a fix soon.

franzpoeschel Oct 7, 2024
Author

Sounds good!
This means that I can use the Kokkos backend as a workaround in the meanwhile?

anagainaru Oct 7, 2024
Maintainer

So, annoyingly the Kokkos backend works fine and the CUDA backend shows zeros.

So for now let me talk about the changes I made to your example over your branch (I will push the changes in 2 min).
You allocate a width (or num columns) of 10 and height (or num rows) of 20 elements.

    cudaMallocPitch(&devPtr, &pitch, width * sizeof(value_type), height);

pitch / sizeof(value_type) is equal to 64 which means you will have an array of 64 columns and 20 rows (with the last 54 elements on each row just padding).

Values are populated on the GPU with consecutive values per row, which means the raw data on the GPU has 64 columns and 20 rows. If I give this to ADIOS2, the raw gpu array will look like:

$ ./bin/bpls MemSelection.bp/ -d gpu_pointer_without_mem_selection -n 64
  int64_t  gpu_pointer_without_mem_selection  {20, 64}
    ( 0, 0)    0 1 2 3 4 5 6 7 8 9 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 1, 0)    10 11 12 13 14 15 16 17 18 19 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 2, 0)    20 21 22 23 24 25 26 27 28 29 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 3, 0)    30 31 32 33 34 35 36 37 38 39 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 4, 0)    40 41 42 43 44 45 46 47 48 49 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 5, 0)    50 51 52 53 54 55 56 57 58 59 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 6, 0)    60 61 62 63 64 65 66 67 68 69 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 7, 0)    70 71 72 73 74 75 76 77 78 79 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 8, 0)    80 81 82 83 84 85 86 87 88 89 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    ( 9, 0)    90 91 92 93 94 95 96 97 98 99 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (10, 0)    100 101 102 103 104 105 106 107 108 109 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (11, 0)    110 111 112 113 114 115 116 117 118 119 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (12, 0)    120 121 122 123 124 125 126 127 128 129 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (13, 0)    130 131 132 133 134 135 136 137 138 139 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (14, 0)    140 141 142 143 144 145 146 147 148 149 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (15, 0)    150 151 152 153 154 155 156 157 158 159 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (16, 0)    160 161 162 163 164 165 166 167 168 169 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (17, 0)    170 171 172 173 174 175 176 177 178 179 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (18, 0)    180 181 182 183 184 185 186 187 188 189 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
    (19, 0)    190 191 192 193 194 195 196 197 198 199 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0

In order to write just the selection ADIOS2 will need to create a variable to store 10 columns and 20 rows (and create the memory selection for 64 columns and 20 rows). This is the values written with the Kokkos backend after my changes to your example:

$ ./bin/bpls MemSelection.bp/ -d gpu_pointer_with_mem_selection -n 10
  int64_t  gpu_pointer_with_mem_selection     {20, 10}
    ( 0,0)    0 1 2 3 4 5 6 7 8 9
    ( 1,0)    10 11 12 13 14 15 16 17 18 19
    ( 2,0)    20 21 22 23 24 25 26 27 28 29
    ( 3,0)    30 31 32 33 34 35 36 37 38 39
    ( 4,0)    40 41 42 43 44 45 46 47 48 49
    ( 5,0)    50 51 52 53 54 55 56 57 58 59
    ( 6,0)    60 61 62 63 64 65 66 67 68 69
    ( 7,0)    70 71 72 73 74 75 76 77 78 79
    ( 8,0)    80 81 82 83 84 85 86 87 88 89
    ( 9,0)    90 91 92 93 94 95 96 97 98 99
    (10,0)    100 101 102 103 104 105 106 107 108 109
    (11,0)    110 111 112 113 114 115 116 117 118 119
    (12,0)    120 121 122 123 124 125 126 127 128 129
    (13,0)    130 131 132 133 134 135 136 137 138 139
    (14,0)    140 141 142 143 144 145 146 147 148 149
    (15,0)    150 151 152 153 154 155 156 157 158 159
    (16,0)    160 161 162 163 164 165 166 167 168 169
    (17,0)    170 171 172 173 174 175 176 177 178 179
    (18,0)    180 181 182 183 184 185 186 187 188 189
    (19,0)    190 191 192 193 194 195 196 197 198 199

Since you use cudaMemcpy2D to move the data, you define the num columns and rows on the CPU the same as the GPU so you don't need to reverse the dimensions.

$ ./bin/bpls MemSelection.bp/ -d cpu_pointer_without_padding -n 10
  int64_t  cpu_pointer_without_padding        {20, 10}
    (0, 0)    0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19
    (1, 0)    20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39
    (2, 0)    40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59
    (3, 0)    60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79
    (4, 0)    80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99
    (5, 0)    100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119
    (6, 0)    120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139
    (7, 0)    140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159
    (8, 0)    160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179
    (9, 0)    180 181 182 183 184 185 186 187 188 189 190 191 192 193 194 195 196 197 198 199

Now, for the CUDA backend we use cudaMemcpy and not cudaMemcpy2D which for some reason returns only zeros (I need to understand why this happens and why KokkosDeepCpy doesn't have this issue). Please try the Kokkos backend, and in the mean time I'll debug cudaMemcpy

anagainaru Oct 7, 2024
Maintainer

For some reason I cannot push on your branch so I created my own with the changes: https://github.com/anagainaru/ADIOS2/tree/franz-testing

Yes, the Kokkos backend should work

franzpoeschel Oct 10, 2024
Author

Thank you for the explanation, this makes sense.
After rebasing on #4359, the example now works as expected.

For some reason I cannot push on your branch

I think, this would have worked only if I had opened a PR here with that branch

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Memory Selection for GPU pointers #4337

{{title}}

Replies: 1 comment 17 replies

{{title}}

{{title}}

{{title}}

{{title}}

{{editor}}'s edit

{{editor}}'s edit

{{title}}

{{editor}}'s edit

{{editor}}'s edit

{{title}}

Select a reply

Memory Selection for GPU pointers #4337

franzpoeschel Sep 9, 2024

Replies: 1 comment · 17 replies

anagainaru Sep 9, 2024 Maintainer

anagainaru Oct 7, 2024 Maintainer

franzpoeschel Oct 7, 2024 Author

anagainaru Oct 7, 2024 Maintainer

anagainaru Oct 7, 2024 Maintainer

franzpoeschel Oct 10, 2024 Author

franzpoeschel
Sep 9, 2024

Replies: 1 comment 17 replies

anagainaru
Sep 9, 2024
Maintainer

anagainaru Oct 7, 2024
Maintainer

franzpoeschel Oct 7, 2024
Author

anagainaru Oct 7, 2024
Maintainer

anagainaru Oct 7, 2024
Maintainer

franzpoeschel Oct 10, 2024
Author