launch.h

#pragma once
#include "ff/energy.h"
#include "tool/cudalib.h"
#include "tool/gpucard.h"
 
namespace tinker {
 
template <class K, class... Ts>
void launch_k3s(cudaStream_t st, size_t sh, int bs, int np, K k, Ts&&... a)
{
   int gs = (np + bs - 1) / bs;
   k<<<gs, bs, sh, st>>>(std::forward<Ts>(a)...);
}
 
template <class K, class... Ts>
void launch_k3b(cudaStream_t st, size_t sh, int bs, int np, K k, Ts&&... a)
{
   np = std::min(np, nelem_buffer);
   int gs = (np + bs - 1) / bs;
   k<<<gs, bs, sh, st>>>(std::forward<Ts>(a)...);
}
 
template <class K, class... Ts>
void launch_k2s(cudaStream_t st, int bs, int np, K k, Ts&&... a)
{
   const size_t sh = 0;
   launch_k3s(st, sh, bs, np, k, std::forward<Ts>(a)...);
}
 
template <class K, class... Ts>
void launch_k2b(cudaStream_t st, int bs, int np, K k, Ts&&... a)
{
   const size_t sh = 0;
   launch_k3b(st, sh, bs, np, k, std::forward<Ts>(a)...);
}
 
template <class K, class... Ts>
void launch_k1s(cudaStream_t st, int np, K k, Ts&&... a)
{
   const int bs = BLOCK_DIM;
   launch_k2s(st, bs, np, k, std::forward<Ts>(a)...);
}
 
template <class K, class... Ts>
void launch_k1b(cudaStream_t st, int np, K k, Ts&&... a)
{
   const int bs = BLOCK_DIM;
   launch_k2b(st, bs, np, k, std::forward<Ts>(a)...);
}
 
#define ITHREAD threadIdx.x + blockIdx.x* blockDim.x
#define STRIDE  blockDim.x* gridDim.x
 
}