cuda_utilities.h

 
#pragma once
 
#include <string>
 
// Local Includes
#include "../global/global.h"
#include "../global/global_cuda.h"
#include "../utils/gpu.hpp"
 
namespace cuda_utilities
{
inline __host__ __device__ void compute3DIndices(int const &id, int const &nx, int const &ny, int &xid, int &yid,
                                                 int &zid)
{
  zid = id / (nx * ny);
  yid = (id - zid * nx * ny) / nx;
  xid = id - zid * nx * ny - yid * nx;
}
 
inline __host__ __device__ int compute1DIndex(int const &xid, int const &yid, int const &zid, int const &nx,
                                              int const &ny)
{
  return xid + yid * nx + zid * nx * ny;
}
 
inline __host__ __device__ void Get_Real_Indices(int const &n_ghost, int const &nx, int const &ny, int const &nz,
                                                 int &is, int &ie, int &js, int &je, int &ks, int &ke)
{
  is = n_ghost;
  ie = nx - n_ghost;
  if (ny == 1) {
    js = 0;
    je = 1;
  } else {
    js = n_ghost;
    je = ny - n_ghost;
  }
  if (nz == 1) {
    ks = 0;
    ke = 1;
  } else {
    ks = n_ghost;
    ke = nz - n_ghost;
  }
}
 
inline void initGpuMemory(Real *ptr, size_t N) { GPU_Error_Check(cudaMemset(ptr, 0, N)); }
 
// =====================================================================
template <typename T>
struct AutomaticLaunchParams {
 public:
  AutomaticLaunchParams(T &kernel, size_t numElements = 0)
  {
    // Get the max number of threads per block allowed by this kernel
    cudaFuncAttributes kernel_attrs{};
    GPU_Error_Check(cudaFuncGetAttributes(&kernel_attrs, reinterpret_cast<const void *>(&kernel)));
 
    // Determine the launch parameters
    cudaOccupancyMaxPotentialBlockSize(&numBlocks, &threadsPerBlock, kernel, 0, kernel_attrs.maxThreadsPerBlock);
 
    if (numElements > 0) {
      // This line is needed to check that threadsPerBlock isn't zero. Somewhere inside
      // cudaOccupancyMaxPotentialBlockSize threadsPerBlock can be zero according to clang-tidy so this line sets it to
      // a more reasonable value
      threadsPerBlock = (threadsPerBlock == 0) ? TPB : threadsPerBlock;
 
      // Compute the number of blocks
      numBlocks = (numElements + threadsPerBlock - 1) / threadsPerBlock;
    }
  }
 
  ~AutomaticLaunchParams() = default;
 
  int get_threadsPerBlock() const { return threadsPerBlock; }
  int get_numBlocks() const { return numBlocks; }
 
 private:
  int threadsPerBlock;
  int numBlocks;
};
// =====================================================================
 
// =====================================================================
void Print_GPU_Memory_Usage(std::string const &additional_text = "");
// =====================================================================
}  // end namespace cuda_utilities