79 return omp_get_thread_num();
86 return omp_get_max_threads();
117 MPI_Comm amps_CommWorld = MPI_COMM_NULL;
118 MPI_Comm amps_CommNode = MPI_COMM_NULL;
120 MPI_Comm_dup(MPI_COMM_WORLD, &s_CommWorld);
121 MPI_Comm_size(amps_CommWorld, &s_size);
122 MPI_Comm_rank(amps_CommWorld, &s_rank);
126 MPI_Comm_split_type(amps_CommWorld, MPI_COMM_TYPE_SHARED, 0, MPI_INFO_NULL, &s_CommNode);
129 char processor_name[MPI_MAX_PROCESSOR_NAME];
131 MPI_Get_processor_name(processor_name, &namelen);
132 uint32_t checkSum = Adler32((
unsigned char*)processor_name, namelen);
136 MPI_Comm_split(amps_CommWorld, checkSum, amps_rank, &s_CommNode);
138 MPI_Comm_rank(amps_CommNode, &s_node_rank);
139 MPI_Comm_size(amps_CommNode, &s_node_size);
144 if (deviceCount > 1) {
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;
152 std::cout <<
"(Node 0) MPI ranks see single GPU device each." << std::endl;
167 #if defined(USE_GPU) && defined(__CUDACC__)
170 #if defined(USE_GPU) && defined(__HIP_PLATFORM_HCC___)
176 #if defined(USE_GPU) && defined(__CUDACC__)
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");
197 int *leastPriority =
new int;
198 int *greatestPriority =
new int;
200 if (*leastPriority==*greatestPriority) {
201 printf(
"Warning when initializing GPU streams: minimum and maximum stream priority are identical! %d == %d \n",*leastPriority, *greatestPriority);
203 for (uint
i=0;
i<maxNThreads; ++
i) {
236 for (uint
i=0;
i<maxNThreads; ++
i) {
267 const size_t local_cells_size,
const size_t ghost_cells_size) {
274 sizeof(std::array<vmesh::MeshParameters,MAX_VMESH_PARAMETERS_COUNT>)
278 size_t vlasovBuffers = 0;
279 size_t batchBuffers = 0;
281 size_t accBuffers = 0;
288 size_t transBuffers = 0;
290 transBuffers +=
sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>);
291 transBuffers +=
unionOfBlocksSet->bucket_count() *
sizeof(Hashinator::hash_pair<vmesh::GlobalID,vmesh::LocalID>);
294 transBuffers +=
sizeof(split::SplitVector<vmesh::GlobalID>);
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);
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;
324 logFile<<
" Total: "<<sum_mb<<
" Mbytes"<<std::endl;
325 logFile<<
" Reported Hardware use: "<<used_mb<<
" Mbytes"<<std::endl;
326 logFile<<
" =================================="<<std::endl;
336 const uint maxBlockCount
360 const uint maxBlockCount
370 std::array<uint, 3> s = {c0,c1,c2};
371 std::sort(s.begin(), s.end());
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)));
376 size_t probeCubeExtentsFlat = s[1]*s[2];
377 probeCubeExtentsFlat = 2*Hashinator::defaults::MAX_BLOCKSIZE * (1 + ((probeCubeExtentsFlat - 1) / (2*Hashinator::defaults::MAX_BLOCKSIZE)));
396 probeAllocation = (1 + ((probeAllocation - 1) / (
WID3 *
sizeof(
Realf)))) * (
WID3 *
sizeof(
Realf));
411 uint smallestAllocation = std::numeric_limits<uint>::max();
418 return smallestAllocation;
423 uint blockAllocationCount
434 size_t blockDataAllocation = blockAllocationCount *
WID3 *
sizeof(
Realf);
449 blockDataAllocation = (1 + ((blockDataAllocation - 1) / (
WID3 *
sizeof(
Realf)))) * (
WID3 *
sizeof(
Realf));
569 uint firstAllocationCount,
571 uint columnSetAllocationCount
573 uint columnAllocationCount;
574 if (columnSetAllocationCount==0) {
585 columnSetAllocationCount = columnAllocationCount;
587 columnAllocationCount = firstAllocationCount;
619 if (largestVmesh > 0) {
620 const vmesh::LocalID HashmapReqSize = ceil(log2((
int)largestVmesh)) +2;
623 void *buf0 = malloc(
sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>));
624 unionOfBlocksSet = ::new (buf0) Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>(HashmapReqSize);
631 void *buf0 = malloc(
sizeof(Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>));
632 unionOfBlocksSet = ::new (buf0) Hashinator::Hashmap<vmesh::GlobalID,vmesh::LocalID>(HashmapReqSize);
642 if (unionSetSize > 0) {
645 void *buf0 = malloc(
sizeof(split::SplitVector<vmesh::GlobalID>));
646 unionOfBlocks = ::new (buf0) split::SplitVector<vmesh::GlobalID>(unionSetSize);
#define gpuDeviceGetStreamPriorityRange
#define gpuDevAttrMaxBlocksPerMultiprocessor
cudaStream_t gpuStreamList[]
#define gpuStreamSynchronize
#define gpuMemcpyHostToDevice
#define gpuStreamCreateWithPriority
#define gpuGetDeviceProperties
#define gpuDeviceSynchronize
#define gpuGetDeviceCount
#define gpuDeviceGetAttribute
__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)
__host__ void gpu_acc_allocate(uint maxBlockCount)
__host__ void gpu_clear_device()
__host__ void gpu_acc_allocate_perthread(uint allocID, uint firstAllocationCount, uint columnSetAllocationCount)
GPUMemoryManager gpuMemoryManager
__host__ void gpu_vlasov_allocate_perthread(uint allocID, uint blockAllocationCount)
std::vector< uint > gpu_vlasov_allocatedSize
split::SplitVector< vmesh::GlobalID > * dev_unionOfBlocks
__host__ void gpu_init_device()
__host__ int gpu_getDevice()
ColumnOffsets * host_columnOffsetData
Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > * dev_unionOfBlocksSet
split::SplitVector< vmesh::GlobalID > * unionOfBlocks
gpuStream_t gpuPriorityStreamList[MAXCPUTHREADS]
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)
__host__ void gpu_trans_allocate(cuint nAllCells, cuint largestVmesh, cuint unionSetSize)
__host__ gpuStream_t gpu_getPriorityStream()
uint gpu_allocated_unionSetSize
Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > * unionOfBlocksSet
__host__ void gpu_vlasov_deallocate()
__host__ void gpu_acc_deallocate()
uint gpu_allocated_largestVmeshSizePower
__host__ gpuStream_t gpu_getStream()
__host__ void gpu_batch_allocate(uint nCells, uint maxNeighbours)
__host__ uint gpu_vlasov_getSmallestAllocation()
size_t gpu_probeFlattenedSize
__host__ void gpu_calculateProbeAllocation(const uint maxBlockCount)
uint gpu_largest_columnCount
__host__ uint gpu_getThread()
__host__ uint gpu_getMaxThreads()
__host__ uint gpu_getAllocationCount()
std::vector< size_t > gpu_vlasov_subPointers
__host__ void gpu_trans_deallocate()
__host__ void gpu_vlasov_allocate(const uint maxBlockCount)
int gpuMultiProcessorCount
static const uint VLASOV_BUFFER_MINBLOCKS
static const int GPU_PROBEFLAT_N
#define HOST_ALLOCATE_WITH_BUFFER(object, member, bytes, buffer)
static const double BLOCK_ALLOCATION_PADDING
#define ALLOCATE_WITH_BUFFER(object, member, bytes, buffer)
#define HOST_ALLOCATE_GPU(object, member, bytes)
#define SUBPOINTER_HOST_ALLOCATE(object, member, index, bytes)
#define SUBPOINTER_ALLOCATE(object, member, index, bytes)
#define CREATE_SUBPOINTERS(object, member, amount)
static const uint VLASOV_BUFFER_MINCOLUMNS
#define GET_POINTER(object, type, member)
static const double BLOCK_ALLOCATION_FACTOR
#define CREATE_UNIQUE_POINTER(object, member)
#define ALLOCATE_GPU(object, member, bytes)
#define SET_SUBPOINTER(object, type, member, index, subPointerIndex)
__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()
ARCH_HOSTDEV MeshWrapper * getMeshWrapper()
std::vector< species::Species > particleSpecies
static uint GPUallocations
std::array< vmesh::MeshParameters, MAX_VMESH_PARAMETERS_COUNT > * velocityMeshes
static ARCH_HOSTDEV VecSimple< T > max(VecSimple< T > const &l, VecSimple< T > const &r)