Method initialize_gpu

Src/Base/AMReX_GpuDevice.cpp:493–689

}

void
Device::initialize_gpu (bool minimal)
{
    amrex::ignore_unused(minimal);

#ifdef AMREX_USE_GPU

    if (gpu_stream_pool.size() != max_gpu_streams) {
        // no copy/move constructor for std::mutex
        gpu_stream_pool = Vector<StreamManager>(max_gpu_streams);
    }

#ifdef AMREX_USE_HIP

    AMREX_HIP_SAFE_CALL(hipGetDeviceProperties(&device_prop, device_id));

    AMREX_ALWAYS_ASSERT_WITH_MESSAGE(warp_size == device_prop.warpSize, "Incorrect warp size");

    // check compute capability

    // AMD devices do not support shared cache banking.

    for (int i = 0; i < max_gpu_streams; ++i) {
        AMREX_HIP_SAFE_CALL(hipStreamCreate(&gpu_stream_pool[i].getStream()));
    }

#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
    AMREX_HIP_SAFE_CALL(hipDeviceGetAttribute(&memory_pools_supported, hipDeviceAttributeMemoryPoolsSupported, device_id));
#endif

#elif defined(AMREX_USE_CUDA)
    AMREX_CUDA_SAFE_CALL(cudaGetDeviceProperties(&device_prop, device_id));

    AMREX_ALWAYS_ASSERT_WITH_MESSAGE(device_prop.major >= 4 || (device_prop.major == 3 && device_prop.minor >= 5),
                                     "Compute capability must be >= 3.5");

#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
    cudaDeviceGetAttribute(&memory_pools_supported, cudaDevAttrMemoryPoolsSupported, device_id);
#endif

#if defined(CUDART_VERSION) && (CUDART_VERSION < 12040)
    if ( ! minimal ) {
        if (sizeof(Real) == 8) {
            AMREX_CUDA_SAFE_CALL(cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte));
        } else if (sizeof(Real) == 4) {
            AMREX_CUDA_SAFE_CALL(cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeFourByte));
        }
    }
#endif

    for (int i = 0; i < max_gpu_streams; ++i) {
        AMREX_CUDA_SAFE_CALL(cudaStreamCreate(&gpu_stream_pool[i].getStream()));
#ifdef AMREX_USE_ACC
        acc_set_cuda_stream(i, gpu_stream_pool[i].getStream());
#endif
    }

    AMREX_ALWAYS_ASSERT_WITH_MESSAGE(warp_size == device_prop.warpSize, "Incorrect warp size");