// Copyright (c) 2010-2023, Lawrence Livermore National Security, LLC. Produced // at the Lawrence Livermore National Laboratory. All Rights reserved. See files // LICENSE and NOTICE for details. LLNL-CODE-806117. // // This file is part of the MFEM library. For more information and source code // availability visit https://mfem.org. // // MFEM is free software; you can redistribute it and/or modify it under the // terms of the BSD-3 license. We welcome feedback and contributions, see file // CONTRIBUTING.md for details. #ifndef MFEM_CUDA_HPP #define MFEM_CUDA_HPP #include "../config/config.hpp" #include "error.hpp" // CUDA block size used by MFEM. #define MFEM_CUDA_BLOCKS 256 #ifdef MFEM_USE_CUDA #define MFEM_USE_CUDA_OR_HIP #define MFEM_DEVICE __device__ #define MFEM_LAMBDA __host__ #define MFEM_HOST_DEVICE __host__ __device__ #define MFEM_DEVICE_SYNC MFEM_GPU_CHECK(cudaDeviceSynchronize()) #define MFEM_STREAM_SYNC MFEM_GPU_CHECK(cudaStreamSynchronize(0)) // Define a CUDA error check macro, MFEM_GPU_CHECK(x), where x returns/is of // type 'cudaError_t'. This macro evaluates 'x' and raises an error if the // result is not cudaSuccess. #define MFEM_GPU_CHECK(x) \ do \ { \ cudaError_t err = (x); \ if (err != cudaSuccess) \ { \ mfem_cuda_error(err, #x, _MFEM_FUNC_NAME, __FILE__, __LINE__); \ } \ } \ while (0) #endif // MFEM_USE_CUDA // Define the MFEM inner threading macros #if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__) #define MFEM_SHARED __shared__ #define MFEM_SYNC_THREAD __syncthreads() #define MFEM_BLOCK_ID(k) blockIdx.k #define MFEM_THREAD_ID(k) threadIdx.k #define MFEM_THREAD_SIZE(k) blockDim.k #define MFEM_FOREACH_THREAD(i,k,N) for(int i=threadIdx.k; i