Vlasiator ebf0dd394 on dev (v5.4.0 + 1054 commits)
Loading...
Searching...
No Matches
gpu_base.cpp
Go to the documentation of this file.
1/*
2 * This file is part of Vlasiator.
3 * Copyright 2010-2025 Finnish Meteorological Institute and University of Helsinki
4 *
5 * For details of usage, see the COPYING file and read the "Rules of the Road"
6 * at http://www.physics.helsinki.fi/vlasiator/
7 *
8 * This program is free software; you can redistribute it and/or modify
9 * it under the terms of the GNU General Public License as published by
10 * the Free Software Foundation; either version 2 of the License, or
11 * (at your option) any later version.
12 *
13 * This program is distributed in the hope that it will be useful,
14 * but WITHOUT ANY WARRANTY; without even the implied warranty of
15 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
16 * GNU General Public License for more details.
17 *
18 * You should have received a copy of the GNU General Public License along
19 * with this program; if not, write to the Free Software Foundation, Inc.,
20 * 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
21 */
22
23#include <stdio.h>
24#include <iostream>
25#include "common.h"
26#include "mpi.h"
27
28#include "gpu_base.hpp"
30#include "object_wrapper.h"
36
37#include "logger.h"
38
39// #define MAXCPUTHREADS 64 now in gpu_base.hpp
40
41// Device properties
45
46extern Logger logFile;
49
50// Allocate pointers for per-thread memory regions
53
54// Pointers to buffers used in acceleration
56// Counts used in acceleration
58
59// Buffers, Vector and set for use in translation
60split::SplitVector<vmesh::GlobalID> *unionOfBlocks=NULL, *dev_unionOfBlocks=NULL;
61Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID> *unionOfBlocksSet=NULL, *dev_unionOfBlocksSet=NULL;
62
63// Memory manager
65
66// Counter for how many parallel vlasov buffers are allocated
68// Counter for how large each allocation is
69std::vector<uint> gpu_vlasov_allocatedSize;
70std::vector<size_t> gpu_vlasov_subPointers;
71
72// counters for allocated sizes in translation
76
77__host__ uint gpu_getThread() {
78#ifdef _OPENMP
79 return omp_get_thread_num();
80#else
81 return 0;
82#endif
83}
84__host__ uint gpu_getMaxThreads() {
85#ifdef _OPENMP
86 return omp_get_max_threads();
87#else
88 return 1;
89#endif
90}
91
92unsigned int nextPowerOfTwo(unsigned int n) {
93 if (n == 0) return 1;
94 n--; // Handle exact powers of two
95 n |= n >> 1;
96 n |= n >> 2;
97 n |= n >> 4;
98 n |= n >> 8;
99 n |= n >> 16;
100 return n + 1;
101}
102
103__host__ void gpu_init_device() {
104 const uint maxNThreads = gpu_getMaxThreads();
105 int deviceCount;
106 // CHK_ERR( gpuFree(0));
107 CHK_ERR( gpuGetDeviceCount(&deviceCount) );
108 //printf("GPU device count %d with %d threads/streams\n",deviceCount,maxNThreads);
109
110 /* Create communicator with one rank per compute node to identify which GPU to use */
111 int amps_size;
112 int amps_rank;
113 int amps_node_rank;
114 int amps_node_size;
115 // int amps_write_rank;
116 // int amps_write_size;
117 MPI_Comm amps_CommWorld = MPI_COMM_NULL;
118 MPI_Comm amps_CommNode = MPI_COMM_NULL;
119
120 MPI_Comm_dup(MPI_COMM_WORLD, &amps_CommWorld);
121 MPI_Comm_size(amps_CommWorld, &amps_size);
122 MPI_Comm_rank(amps_CommWorld, &amps_rank);
123
124 /* Create communicator with one rank per compute node */
125#if MPI_VERSION >= 3
126 MPI_Comm_split_type(amps_CommWorld, MPI_COMM_TYPE_SHARED, 0, MPI_INFO_NULL, &amps_CommNode);
127#else
128 /* Split the node level communicator based on Adler32 hash keys of processor name */
129 char processor_name[MPI_MAX_PROCESSOR_NAME];
130 int namelen;
131 MPI_Get_processor_name(processor_name, &namelen);
132 uint32_t checkSum = Adler32((unsigned char*)processor_name, namelen);
133 /* Comm split only accepts non-negative numbers */
134 /* Not super great for hashing purposes but hoping MPI-3 code will be used on most cases */
135 checkSum &= INT_MAX;
136 MPI_Comm_split(amps_CommWorld, checkSum, amps_rank, &amps_CommNode);
137#endif
138 MPI_Comm_rank(amps_CommNode, &amps_node_rank);
139 MPI_Comm_size(amps_CommNode, &amps_node_size);
140 myRank = amps_rank;
141
142 // if only one visible device, assume MPI system handles device visibility and just use the only visible one.
143 if (amps_rank == MASTER_RANK) {
144 if (deviceCount > 1) {
145 // If more than one device is visible, issue warning to user along with suggestion to use SLURM options.
146 std::cout << "(Node 0) WARNING! MPI ranks see "<<deviceCount<<" GPU devices each." << std::endl;
147 std::cout << " Recommended usage is to utilize SLURM for showing only single GPU device per MPI rank:" << std::endl;
148 std::cout << " export CUDA_VISIBLE_DEVICES=\\$SLURM_LOCALID" << std::endl;
149 std::cout << " or" << std::endl;
150 std::cout << " export ROCR_VISIBLE_DEVICES=\\$SLURM_LOCALID" << std::endl;
151 } else {
152 std::cout << "(Node 0) MPI ranks see single GPU device each." << std::endl;
153 }
154 }
157
158 // Decide on number of allocations to prepare
159 const uint nBaseCells = P::xcells_ini * P::ycells_ini * P::zcells_ini;
160 allocationCount = (nBaseCells == 1) ? 1 : P::GPUallocations;
161
162 // Get device properties
163 gpuDeviceProp prop;
165 gpuMultiProcessorCount = prop.multiProcessorCount;
166 threadsPerMP = prop.maxThreadsPerMultiProcessor;
167 #if defined(USE_GPU) && defined(__CUDACC__)
169 #endif
170 #if defined(USE_GPU) && defined(__HIP_PLATFORM_HCC___)
171 blocksPerMP = threadsPerMP/GPUTHREADS; // This should be the maximum number of wavefronts per CU
172 #endif
173
174
175 // Query device capabilities (only for CUDA, not needed for HIP)
176 #if defined(USE_GPU) && defined(__CUDACC__)
177 int supportedMode;
178 CHK_ERR( cudaDeviceGetAttribute (&supportedMode, cudaDevAttrConcurrentManagedAccess, myDevice) );
179 if (supportedMode==0) {
180 printf("Error! Current GPU device does not support concurrent managed memory access from several streams.\n");
181 printf("Please switch to a more recent CUDA compute architecture.\n");
182 abort();
183 }
184 #endif
185
186 // Pre-generate streams, allocate return pointers
187 CREATE_SUBPOINTERS(gpuMemoryManager, returnReal, maxNThreads);
188 CREATE_SUBPOINTERS(gpuMemoryManager, returnRealf, maxNThreads);
189 CREATE_SUBPOINTERS(gpuMemoryManager, returnLID, maxNThreads);
190 CREATE_SUBPOINTERS(gpuMemoryManager, host_returnReal, maxNThreads);
191 CREATE_SUBPOINTERS(gpuMemoryManager, host_returnRealf, maxNThreads);
192 CREATE_SUBPOINTERS(gpuMemoryManager, host_returnLID, maxNThreads);
193 CREATE_SUBPOINTERS(gpuMemoryManager, gpuInitBuffer, maxNThreads);
194 CREATE_SUBPOINTERS(gpuMemoryManager, gpuInitBlocks, maxNThreads);
195
196
197 int *leastPriority = new int; // likely 0
198 int *greatestPriority = new int; // likely -1
199 CHK_ERR( gpuDeviceGetStreamPriorityRange (leastPriority, greatestPriority) );
200 if (*leastPriority==*greatestPriority) {
201 printf("Warning when initializing GPU streams: minimum and maximum stream priority are identical! %d == %d \n",*leastPriority, *greatestPriority);
202 }
203 for (uint i=0; i<maxNThreads; ++i) {
206
207 SUBPOINTER_HOST_ALLOCATE(gpuMemoryManager, host_returnReal, i, 8*sizeof(Real));
208 SUBPOINTER_HOST_ALLOCATE(gpuMemoryManager, host_returnRealf, i, 8*sizeof(Realf));
209 SUBPOINTER_HOST_ALLOCATE(gpuMemoryManager, host_returnLID, i, 8*sizeof(vmesh::LocalID));
210
211 SUBPOINTER_ALLOCATE(gpuMemoryManager, returnReal, i, 8*sizeof(Real));
212 SUBPOINTER_ALLOCATE(gpuMemoryManager, returnRealf, i, 8*sizeof(Realf));
214 }
215
216 CREATE_UNIQUE_POINTER(gpuMemoryManager, gpu_cell_indices_to_id);
218 CREATE_UNIQUE_POINTER(gpuMemoryManager, gpu_block_indices_to_probe);
219 CREATE_UNIQUE_POINTER(gpuMemoryManager, my_test_pointer);
220
221 ALLOCATE_GPU(gpuMemoryManager, gpu_cell_indices_to_id, 3*sizeof(uint));
223 ALLOCATE_GPU(gpuMemoryManager, gpu_block_indices_to_probe, 3*sizeof(uint));
225
226 // Using just a single context for whole MPI task
227}
228
229__host__ void gpu_clear_device() {
230 // Deallocate temporary buffers
234 // Destroy streams
235 const uint maxNThreads = gpu_getMaxThreads();
236 for (uint i=0; i<maxNThreads; ++i) {
239 }
241 gpuMemoryManager.freeAll();
242}
243
246}
247
249 const uint thread_id = gpu_getThread();
250 return gpuPriorityStreamList[thread_id];
251}
252
253__host__ int gpu_getDevice() {
254 int device;
255 CHK_ERR( gpuGetDevice(&device) );
256 return device;
257}
258
259__host__ uint gpu_getAllocationCount() {
260 return allocationCount;
261}
262
263/*
264 Memory reporting function
265*/
266int gpu_reportMemory(const size_t local_cells_capacity, const size_t ghost_cells_capacity,
267 const size_t local_cells_size, const size_t ghost_cells_size) {
268 /* Gather total CPU and GPU buffer sizes. Rank 0 reports details,
269 all ranks return sum.
270 */
271 uint maxNThreads = gpu_getMaxThreads();
272
273 size_t miniBuffers =
274 sizeof(std::array<vmesh::MeshParameters,MAX_VMESH_PARAMETERS_COUNT>) // velocityMeshes_upload
275 + sizeof(vmesh::MeshWrapper); // MWdev
276 // DT reduction buffers are deallocated every step (GPUTODO, make persistent)
277
278 size_t vlasovBuffers = 0;
279 size_t batchBuffers = 0;
280
281 size_t accBuffers = 0;
282 for (uint i=0; i<allocationCount; ++i) {
284 accBuffers += host_columnOffsetData[i].capacityInBytes(); // struct contents
285 }
286 }
287
288 size_t transBuffers = 0;
289 if (unionOfBlocksSet) {
290 transBuffers += sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>);
291 transBuffers += unionOfBlocksSet->bucket_count() * sizeof(Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>);
292 }
293 if (unionOfBlocks) {
294 transBuffers += sizeof(split::SplitVector<vmesh::GlobalID>);
295 transBuffers += unionOfBlocks->capacity() * sizeof(vmesh::GlobalID);
296 }
297 // Remote neighbor contribution buffers are in unified memory but deallocated after each use
298
299 size_t memoryManagerCapacity = gpuMemoryManager.totalGpuAllocation();
300
301 size_t free_byte ;
302 size_t total_byte ;
303 CHK_ERR( gpuMemGetInfo( &free_byte, &total_byte) );
304 size_t used_mb = (total_byte-free_byte)/(1024*1024);
305 size_t sum_mb = (miniBuffers+batchBuffers+vlasovBuffers+accBuffers+transBuffers+local_cells_capacity+ghost_cells_capacity+memoryManagerCapacity)/(1024*1024);
306 size_t local_req_mb = local_cells_size/(1024*1024);
307 size_t ghost_req_mb = ghost_cells_size/(1024*1024);
308
309 if (myRank==0) {
310 logFile<<" =================================="<<std::endl;
311 logFile<<" GPU Memory report"<<std::endl;
312 logFile<<" mini-buffers: "<<miniBuffers/(1024*1024)<<" Mbytes"<<std::endl;
313 logFile<<" Batch buffers: "<<batchBuffers/(1024*1024)<<" Mbytes"<<std::endl;
314 logFile<<" Vlasov buffers: "<<vlasovBuffers/(1024*1024)<<" Mbytes"<<std::endl;
315 logFile<<" Acceleration buffers: "<<accBuffers/(1024*1024)<<" Mbytes"<<std::endl;
316 logFile<<" Translation buffers: "<<transBuffers/(1024*1024)<<" Mbytes"<<std::endl;
317 logFile<<" Local cells: "<<local_cells_capacity/(1024*1024)<<" Mbytes"<<std::endl;
318 logFile<<" Ghost cells: "<<ghost_cells_capacity/(1024*1024)<<" Mbytes"<<std::endl;
319 logFile<<" Memory manager: "<<memoryManagerCapacity/(1024*1024)<<" Mbytes"<<std::endl;
320 if (local_req_mb || ghost_req_mb) {
321 logFile<<" Local cells required: "<<local_req_mb<<" Mbytes"<<std::endl;
322 logFile<<" Ghost cells required: "<<ghost_req_mb<<" Mbytes"<<std::endl;
323 }
324 logFile<<" Total: "<<sum_mb<<" Mbytes"<<std::endl;
325 logFile<<" Reported Hardware use: "<<used_mb<<" Mbytes"<<std::endl;
326 logFile<<" =================================="<<std::endl;
327 }
328 return sum_mb;
329}
330
331/*
332 Top-level GPU memory allocation function.
333 This is called from within non-threaded regions so does not perform async.
334 */
336 const uint maxBlockCount // Largest found vmesh size
337 ) {
338 // Always prepare for at least VLASOV_BUFFER_MINBLOCKS blocks
339 const uint maxBlocksPerCell = max(VLASOV_BUFFER_MINBLOCKS, maxBlockCount);
340
341 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_blockDataOrdered);
344 HOST_ALLOCATE_GPU(gpuMemoryManager, host_blockDataOrdered, allocationCount*sizeof(Realf*));
345
346 // per-buffer allocations
347 for (uint i=0; i<allocationCount; ++i) {
348 gpu_vlasov_allocate_perthread(i, maxBlocksPerCell);
349 }
350
351 // Above function stores buffer pointers in host_blockDataOrdered, copy pointers to dev_blockDataOrdered
353}
354
355/*
356 Top-level GPU memory allocation function.
357 This is called from within non-threaded regions so does not perform async.
358 */
360 const uint maxBlockCount // Largest found vmesh size
361 ) {
362 // Always prepare for at least VLASOV_BUFFER_MINBLOCKS blocks
363 const uint maxBlocksPerCell = max(VLASOV_BUFFER_MINBLOCKS, maxBlockCount);
364
365 // Evaluate required size for acceleration probe cube (based on largest population)
366 for (uint popID=0; popID<getObjectWrapper().particleSpecies.size(); ++popID) {
367 const uint c0 = (*vmesh::getMeshWrapper()->velocityMeshes)[popID].gridLength[0];
368 const uint c1 = (*vmesh::getMeshWrapper()->velocityMeshes)[popID].gridLength[1];
369 const uint c2 = (*vmesh::getMeshWrapper()->velocityMeshes)[popID].gridLength[2];
370 std::array<uint, 3> s = {c0,c1,c2};
371 std::sort(s.begin(), s.end());
372 // Round values up to nearest 2*Hashinator::defaults::MAX_BLOCKSIZE
373 size_t probeCubeExtentsFull = s[0]*s[1]*s[2];
374 probeCubeExtentsFull = 2*Hashinator::defaults::MAX_BLOCKSIZE * (1 + ((probeCubeExtentsFull - 1) / (2*Hashinator::defaults::MAX_BLOCKSIZE)));
375 gpu_probeFullSize = std::max(gpu_probeFullSize,probeCubeExtentsFull);
376 size_t probeCubeExtentsFlat = s[1]*s[2];
377 probeCubeExtentsFlat = 2*Hashinator::defaults::MAX_BLOCKSIZE * (1 + ((probeCubeExtentsFlat - 1) / (2*Hashinator::defaults::MAX_BLOCKSIZE)));
378 gpu_probeFlattenedSize = std::max(gpu_probeFlattenedSize,probeCubeExtentsFlat);
379 }
380
382 /*
383 CUDA C Programming Guide
384 6.3.2. Device Memory Accesses (June 2025)
385 "Any address of a variable residing in global memory or returned by one of the memory allocation routines from the driver or
386 runtime API is always aligned to at least 256 bytes."
387
388 ROCm documentation
389 HIP 6.4.43483 Documentation for hipMallocPitch
390 "Currently the alignment is set to 128 bytes"
391
392 Thus, our mallocs should be in increments of 256 bytes. WID3 is at least 64, and len(Realf) is at least 4, so this is true in all
393 cases. Still, let us ensure (just to be sure) that probe cube addressing does not break alignment.
394 And in fact let's use the block memory size as the stride.
395 */
396 probeAllocation = (1 + ((probeAllocation - 1) / (WID3 * sizeof(Realf)))) * (WID3 * sizeof(Realf));
397 gpu_probeStride = max(gpu_probeStride,probeAllocation);
398}
399
400/* Deallocation at end of simulation */
401__host__ void gpu_vlasov_deallocate() {
402 while(gpu_vlasov_allocatedSize.size() < allocationCount){ //Make sure the gpu_vlasov_allocatedSize has enough elements
403 gpu_vlasov_allocatedSize.push_back(0);
404 }
405 for (uint i=0; i<allocationCount; ++i) {
407 }
408}
409
411 uint smallestAllocation = std::numeric_limits<uint>::max();
412 while(gpu_vlasov_allocatedSize.size() < allocationCount){ //Make sure the gpu_vlasov_allocatedSize has enough elements
413 gpu_vlasov_allocatedSize.push_back(0);
414 }
415 for (uint i=0; i<allocationCount; ++i) {
416 smallestAllocation = std::min(smallestAllocation,gpu_vlasov_allocatedSize[i]);
417 }
418 return smallestAllocation;
419}
420
422 uint allocID,
423 uint blockAllocationCount
424 ) {
425 while(gpu_vlasov_allocatedSize.size() < allocationCount){ //Make sure the gpu_vlasov_allocatedSize has enough elements
426 gpu_vlasov_allocatedSize.push_back(0);
427 }
428 while(gpu_vlasov_subPointers.size() < allocationCount){ //Make sure the gpu_vlasov_subPointers has enough elements
429 gpu_vlasov_subPointers.push_back(0);
430 }
431
432 // Dual use of blockDataOrdered: use also for acceleration probe cube and its flattened version.
433 // Calculate required size
434 size_t blockDataAllocation = blockAllocationCount * WID3 * sizeof(Realf);
435 /*
436 CUDA C Programming Guide
437 6.3.2. Device Memory Accesses (June 2025)
438 "Any address of a variable residing in global memory or returned by one of the memory allocation routines from the driver or
439 runtime API is always aligned to at least 256 bytes."
440
441 ROCm documentation
442 HIP 6.4.43483 Documentation for hipMallocPitch
443 "Currently the alignment is set to 128 bytes"
444
445 Thus, our mallocs should be in increments of 256 bytes. WID3 is at least 64, and len(Realf) is at least 4, so this is true in all
446 cases. Still, let us ensure (just to be sure) that probe cube addressing does not break alignment.
447 And in fact let's use the block memory size as the stride.
448 */
449 blockDataAllocation = (1 + ((blockDataAllocation - 1) / (WID3 * sizeof(Realf)))) * (WID3 * sizeof(Realf));
450
451 gpuMemoryManager.createPointer(gpu_vlasov_subPointers[allocID]);
452 bool reAllocated = gpuMemoryManager.allocate(gpu_vlasov_subPointers[allocID], blockDataAllocation);
453 SET_SUBPOINTER(gpuMemoryManager, Realf, host_blockDataOrdered, allocID, gpu_vlasov_subPointers[allocID]);
454
455 // Store size of new allocation (in units blocks)
456 if (reAllocated) {
457 gpu_vlasov_allocatedSize[allocID] = blockDataAllocation / (WID3 * sizeof(Realf));
458 }
459}
460
462__host__ void gpu_batch_allocate(uint nCells, uint maxNeighbours) {
463
468 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_lists_with_replace_new);
469 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_lists_delete);
470 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_lists_to_replace);
471 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_lists_with_replace_old);
474 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_nBlocksToChange);
475 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_resizeSuccess);
476 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_overflownElements);
479 CREATE_UNIQUE_POINTER(gpuMemoryManager, host_intersections);
480
483 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_allMaps, 2*nCells*sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>*), BLOCK_ALLOCATION_FACTOR); // note double size
484 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_vbwcl_vec, nCells*sizeof(split::SplitVector<vmesh::GlobalID>*), BLOCK_ALLOCATION_FACTOR);
485 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_lists_with_replace_new, nCells*sizeof(split::SplitVector<vmesh::GlobalID>*), BLOCK_ALLOCATION_FACTOR);
486 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_lists_delete, nCells*sizeof(split::SplitVector<Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>>*), BLOCK_ALLOCATION_FACTOR);
487 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_lists_to_replace, nCells*sizeof(split::SplitVector<Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>>*), BLOCK_ALLOCATION_FACTOR);
488 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_lists_with_replace_old, nCells*sizeof(split::SplitVector<Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>>*), BLOCK_ALLOCATION_FACTOR);
496 HOST_ALLOCATE_WITH_BUFFER(gpuMemoryManager, host_intersections, nCells*4*sizeof(Realf), BLOCK_ALLOCATION_FACTOR);
497
502 CREATE_UNIQUE_POINTER(gpuMemoryManager, dev_lists_with_replace_new);
503 CREATE_UNIQUE_POINTER(gpuMemoryManager, dev_lists_delete);
504 CREATE_UNIQUE_POINTER(gpuMemoryManager, dev_lists_to_replace);
505 CREATE_UNIQUE_POINTER(gpuMemoryManager, dev_lists_with_replace_old);
508 CREATE_UNIQUE_POINTER(gpuMemoryManager, dev_nBlocksToChange);
514
517 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_allMaps, 2*nCells*sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>*), BLOCK_ALLOCATION_FACTOR);
518 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_vbwcl_vec, nCells*sizeof(split::SplitVector<vmesh::GlobalID>*), BLOCK_ALLOCATION_FACTOR);
519 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_lists_with_replace_new, nCells*sizeof(split::SplitVector<vmesh::GlobalID>*), BLOCK_ALLOCATION_FACTOR);
520 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_lists_delete, nCells*sizeof(split::SplitVector<Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>>*), BLOCK_ALLOCATION_FACTOR);
521 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_lists_to_replace, nCells*sizeof(split::SplitVector<Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>>*), BLOCK_ALLOCATION_FACTOR);
522 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_lists_with_replace_old, nCells*sizeof(split::SplitVector<Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>>*), BLOCK_ALLOCATION_FACTOR);
528 ALLOCATE_WITH_BUFFER(gpuMemoryManager, dev_minValues, nCells*sizeof(Real), BLOCK_ALLOCATION_FACTOR);
531}
532
533/*
534 Top-level GPU memory allocation function for acceleration-specific column data.
535 This is called from within non-threaded regions so does not perform async.
536 */
537__host__ void gpu_acc_allocate(
538 uint maxBlockCount
539 ) {
540 if (host_columnOffsetData == NULL) {
541 // This would be preferable as would use pinned memory but fails on exit
542 void *buf;
543 CHK_ERR( gpuMallocHost((void**)&buf,allocationCount*sizeof(ColumnOffsets)) );
545 }
548 for (uint i=0; i<allocationCount; ++i) {
549 gpu_acc_allocate_perthread(i,maxBlockCount);
550 }
551 // Above function stores buffer pointers in host_blockDataOrdered, copy pointers to dev_blockDataOrdered
553}
554
555/* Deallocation at end of simulation */
556__host__ void gpu_acc_deallocate() {
557 if (host_columnOffsetData != NULL) {
558 // delete[] host_columnOffsetData;
559 }
561}
562
563/*
564 Mid-level GPU memory allocation function for acceleration-specific column data.
565 Supports calling within threaded regions and async operations.
566 */
568 uint allocID,
569 uint firstAllocationCount, // This is treated as maxBlockCount, unless the next
570 //value is nonzero, in which case it is the column allocation count
571 uint columnSetAllocationCount
572 ) {
573 uint columnAllocationCount;
574 if (columnSetAllocationCount==0) {
575 /*
576 Estimate column count from maxblockcount, non-critical if ends up being too small.
577 This makes a rough guess that we have a cubic velocity space domain, and thus one edge
578 of it is the cubic root of the blocks count, and the area is the square of that. Thus,
579 we take the two-thirds power of the block count, and multiply by the padding multiplier
580 to be a bit on the safer side.
581 */
582 columnAllocationCount = BLOCK_ALLOCATION_PADDING * std::pow(firstAllocationCount,0.666);
583 // Ensure a minimum value.
584 columnAllocationCount = std::max(columnAllocationCount,(uint)VLASOV_BUFFER_MINCOLUMNS);
585 columnSetAllocationCount = columnAllocationCount;
586 } else {
587 columnAllocationCount = firstAllocationCount;
588 // Update tracker
589 gpu_largest_columnCount = std::max(columnAllocationCount,gpu_largest_columnCount);
590 }
591
592 // columndata contains several splitvectors. columnData is host/device, but splitvector contents are unified.
593 gpuStream_t stream = gpu_getStream();
594 // Reallocate if necessary
595 if ( (columnAllocationCount > host_columnOffsetData[allocID].capacityCols()) ||
596 (columnSetAllocationCount > host_columnOffsetData[allocID].capacityColSets()) ) {
597 // Also set size to match input
598 host_columnOffsetData[allocID].setSizes(columnAllocationCount*BLOCK_ALLOCATION_PADDING, columnSetAllocationCount*BLOCK_ALLOCATION_PADDING);
600 }
601}
602
603/*
604 Top-level GPU memory allocation function for translation-specific vectors
605 This is called from within non-threaded regions so does not perform async.
606 */
607__host__ void gpu_trans_allocate(
608 cuint nAllCells,
609 cuint largestVmesh,
610 cuint unionSetSize
611 ) {
612 gpuStream_t stream = gpu_getStream();
613 // Vectors with one entry per cell (prefetch to host)
614 if (nAllCells > 0) {
615 // Use batch allocation
616 gpu_batch_allocate(nAllCells);
617 }
618 // Set for collecting union of blocks (prefetched to device)
619 if (largestVmesh > 0) {
620 const vmesh::LocalID HashmapReqSize = ceil(log2((int)largestVmesh)) +2;
622 // New allocation
623 void *buf0 = malloc(sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>));
624 unionOfBlocksSet = ::new (buf0) Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>(HashmapReqSize);
625 dev_unionOfBlocksSet = unionOfBlocksSet->upload<true>(stream); // <true> == optimize to GPU
627 } else {
628 // Ensure allocation
629 if (HashmapReqSize > gpu_allocated_largestVmeshSizePower) {
630 ::delete unionOfBlocksSet;
631 void *buf0 = malloc(sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>));
632 unionOfBlocksSet = ::new (buf0) Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>(HashmapReqSize);
633 dev_unionOfBlocksSet = unionOfBlocksSet->upload<true>(stream); // <true> == optimize to GPU
635 } else {
636 // Ensure map is empty
637 unionOfBlocksSet->clear<false>(Hashinator::targets::device,stream, std::pow(2,gpu_allocated_largestVmeshSizePower));
638 }
639 }
640 }
641 // Vector into which the set contents are read (prefetched to device)
642 if (unionSetSize > 0) {
644 // New allocation
645 void *buf0 = malloc(sizeof(split::SplitVector<vmesh::GlobalID>));
646 unionOfBlocks = ::new (buf0) split::SplitVector<vmesh::GlobalID>(unionSetSize);
647 unionOfBlocks->clear();
648 //unionOfBlocks->optimizeGPU(stream);
649 dev_unionOfBlocks = unionOfBlocks->upload<true>(stream); // <true> == optimize to GPU
650 } else {
651 // Clear is enough
652 unionOfBlocks->clear();
653 unionOfBlocks->reserve(unionSetSize);
654 //unionOfBlocks->optimizeGPU(stream);
655 dev_unionOfBlocks = unionOfBlocks->upload<true>(stream); // <true> == optimize to GPU
656 }
657 gpu_allocated_unionSetSize = unionSetSize;
658 }
659 CHK_ERR( gpuStreamSynchronize(stream) );
660}
661
662/* Deallocation at end of simulation */
663__host__ void gpu_trans_deallocate() {
664 // Deallocate any translation vectors or sets which exist
666 ::delete unionOfBlocksSet;
668 }
670 ::delete unionOfBlocks;
672 }
673}
for i
Definition Dispersion.m:24
#define gpuDeviceGetStreamPriorityRange
#define gpuStream_t
#define gpuDevAttrMaxBlocksPerMultiprocessor
#define gpuGetDevice
cudaStream_t gpuStreamList[]
Definition gpu_base.cpp:51
#define gpuStreamSynchronize
#define gpuMemcpyHostToDevice
#define gpuStreamDestroy
#define CHK_ERR(err)
#define gpuMemcpy
#define gpuStreamCreateWithPriority
#define gpuMemcpyAsync
#define gpuGetDeviceProperties
#define gpuDeviceSynchronize
#define gpuMallocHost
#define gpuDeviceProp
#define gpuMemGetInfo
#define gpuStreamDefault
#define GPUTHREADS
#define gpuGetDeviceCount
#define gpuDeviceGetAttribute
#define MASTER_RANK
Definition common.h:67
const int WID3
Definition common.h:517
const uint32_t cuint
Definition definitions.h:50
float Real
Definition definitions.h:41
float Realf
Definition definitions.h:33
__global__ void vmesh::VelocityMesh **__restrict__ ColumnOffsets * dev_columnOffsetData
__global__ void vmesh::VelocityMesh **__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint *__restrict__ gpu_block_indices_to_id
__global__ void vmesh::VelocityMesh **__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint *__restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID vmesh::LocalID * dev_overflownElements
__global__ void vmesh::VelocityMesh **__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint *__restrict__ const Realf * dev_intersections
__global__ void vmesh::VelocityMesh **__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint *__restrict__ const Realf const int const int const Realf const Realf vmesh::LocalID * dev_resizeSuccess
unsigned int nextPowerOfTwo(unsigned int n)
Definition gpu_base.cpp:92
__host__ void gpu_acc_allocate(uint maxBlockCount)
Definition gpu_base.cpp:537
__host__ void gpu_clear_device()
Definition gpu_base.cpp:229
int blocksPerMP
Definition gpu_base.cpp:43
__host__ void gpu_acc_allocate_perthread(uint allocID, uint firstAllocationCount, uint columnSetAllocationCount)
Definition gpu_base.cpp:567
GPUMemoryManager gpuMemoryManager
Definition gpu_base.cpp:64
__host__ void gpu_vlasov_allocate_perthread(uint allocID, uint blockAllocationCount)
Definition gpu_base.cpp:421
std::vector< uint > gpu_vlasov_allocatedSize
Definition gpu_base.cpp:69
split::SplitVector< vmesh::GlobalID > * dev_unionOfBlocks
Definition gpu_base.cpp:60
__host__ void gpu_init_device()
Definition gpu_base.cpp:103
__host__ int gpu_getDevice()
Definition gpu_base.cpp:253
ColumnOffsets * host_columnOffsetData
Definition gpu_base.cpp:55
Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > * dev_unionOfBlocksSet
Definition gpu_base.cpp:61
split::SplitVector< vmesh::GlobalID > * unionOfBlocks
Definition gpu_base.cpp:60
gpuStream_t gpuPriorityStreamList[MAXCPUTHREADS]
Definition gpu_base.cpp:52
int gpu_reportMemory(const size_t local_cells_capacity, const size_t ghost_cells_capacity, const size_t local_cells_size, const size_t ghost_cells_size)
Definition gpu_base.cpp:266
int myRank
Definition gpu_base.cpp:48
__host__ void gpu_trans_allocate(cuint nAllCells, cuint largestVmesh, cuint unionSetSize)
Definition gpu_base.cpp:607
size_t gpu_probeStride
Definition gpu_base.cpp:57
__host__ gpuStream_t gpu_getPriorityStream()
Definition gpu_base.cpp:248
int threadsPerMP
Definition gpu_base.cpp:44
size_t gpu_probeFullSize
Definition gpu_base.cpp:57
uint gpu_allocated_unionSetSize
Definition gpu_base.cpp:74
Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > * unionOfBlocksSet
Definition gpu_base.cpp:61
Logger logFile
Definition main.cpp:25
__host__ void gpu_vlasov_deallocate()
Definition gpu_base.cpp:401
__host__ void gpu_acc_deallocate()
Definition gpu_base.cpp:556
uint gpu_allocated_largestVmeshSizePower
Definition gpu_base.cpp:73
int myDevice
Definition gpu_base.cpp:47
__host__ gpuStream_t gpu_getStream()
Definition gpu_base.cpp:244
__host__ void gpu_batch_allocate(uint nCells, uint maxNeighbours)
Definition gpu_base.cpp:462
uint allocationCount
Definition gpu_base.cpp:67
__host__ uint gpu_vlasov_getSmallestAllocation()
Definition gpu_base.cpp:410
size_t gpu_probeFlattenedSize
Definition gpu_base.cpp:57
__host__ void gpu_calculateProbeAllocation(const uint maxBlockCount)
Definition gpu_base.cpp:359
uint gpu_largest_columnCount
Definition gpu_base.cpp:75
__host__ uint gpu_getThread()
Definition gpu_base.cpp:77
__host__ uint gpu_getMaxThreads()
Definition gpu_base.cpp:84
__host__ uint gpu_getAllocationCount()
Definition gpu_base.cpp:259
std::vector< size_t > gpu_vlasov_subPointers
Definition gpu_base.cpp:70
__host__ void gpu_trans_deallocate()
Definition gpu_base.cpp:663
__host__ void gpu_vlasov_allocate(const uint maxBlockCount)
Definition gpu_base.cpp:335
int gpuMultiProcessorCount
Definition gpu_base.cpp:42
static const uint VLASOV_BUFFER_MINBLOCKS
Definition gpu_base.hpp:56
#define MAXCPUTHREADS
Definition gpu_base.hpp:73
static const int GPU_PROBEFLAT_N
Definition gpu_base.hpp:65
#define HOST_ALLOCATE_WITH_BUFFER(object, member, bytes, buffer)
Definition gpu_base.hpp:481
static const double BLOCK_ALLOCATION_PADDING
Definition gpu_base.hpp:60
#define ALLOCATE_WITH_BUFFER(object, member, bytes, buffer)
Definition gpu_base.hpp:429
#define HOST_ALLOCATE_GPU(object, member, bytes)
Definition gpu_base.hpp:455
#define SUBPOINTER_HOST_ALLOCATE(object, member, index, bytes)
Definition gpu_base.hpp:534
#define SUBPOINTER_ALLOCATE(object, member, index, bytes)
Definition gpu_base.hpp:507
#define CREATE_SUBPOINTERS(object, member, amount)
Definition gpu_base.hpp:322
static const uint VLASOV_BUFFER_MINCOLUMNS
Definition gpu_base.hpp:57
#define GET_POINTER(object, type, member)
Definition gpu_base.hpp:809
static const double BLOCK_ALLOCATION_FACTOR
Definition gpu_base.hpp:61
#define CREATE_UNIQUE_POINTER(object, member)
Definition gpu_base.hpp:296
#define ALLOCATE_GPU(object, member, bytes)
Definition gpu_base.hpp:403
#define SET_SUBPOINTER(object, type, member, index, subPointerIndex)
Definition gpu_base.hpp:879
__global__ void const Realf const uint *__restrict__ const uint *__restrict__ const vmesh::GlobalID *__restrict__ const uint const uint const uint const Realf const vmesh::VelocityMesh *__restrict__ const vmesh::VelocityBlockContainer Realf Realf ** dev_blockDataOrdered
ObjectWrapper & getObjectWrapper()
Definition main.cpp:33
uint32_t LocalID
Definition definitions.h:60
uint32_t GlobalID
Definition definitions.h:59
ARCH_HOSTDEV MeshWrapper * getMeshWrapper()
std::vector< species::Species > particleSpecies
static uint zcells_ini
Definition parameters.h:50
static uint ycells_ini
Definition parameters.h:49
static uint xcells_ini
Definition parameters.h:48
static uint GPUallocations
Definition parameters.h:79
std::array< vmesh::MeshParameters, MAX_VMESH_PARAMETERS_COUNT > * velocityMeshes
static ARCH_HOSTDEV VecSimple< T > max(VecSimple< T > const &l, VecSimple< T > const &r)