Start using C++ cuda-api-wrapper instead of raw CUDA

1 files changed, 0 insertions, 37 deletions

diff --git a/tangle/LLBM/kernel/executor.h b/tangle/LLBM/kernel/executor.h
index 942918d..cce023b 100644
--- a/tangle/LLBM/kernel/executor.h
+++ b/tangle/LLBM/kernel/executor.h
@@ -30,30 +30,6 @@ __global__ void call_operator(
   }
 }
 
-template <typename OPERATOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
-__global__ void call_operator(
-    LatticeView<DESCRIPTOR,S> lattice
-  , bool* mask
-  , ARGS... args
-) {
-  const std::size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
-  if (!(gid < lattice.cuboid.volume) || !mask[gid]) {
-      return;
-  }
-
-  S f_curr[DESCRIPTOR::q];
-  S f_next[DESCRIPTOR::q];
-  S* preshifted_f[DESCRIPTOR::q];
-  for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
-    preshifted_f[iPop] = lattice.pop(iPop, gid);
-    f_curr[iPop] = *preshifted_f[iPop];
-  }
-  OPERATOR::template apply<T,S>(DESCRIPTOR(), f_curr, f_next, gid, std::forward<ARGS>(args)...);
-  for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
-    *preshifted_f[iPop] = f_next[iPop];
-  }
-}
-
 template <typename FUNCTOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
 __global__ void call_functor(
     LatticeView<DESCRIPTOR,S> lattice
@@ -110,19 +86,6 @@ __global__ void call_operator_using_list(
   OPERATOR::template apply<T,S>(lattice, index, count, std::forward<ARGS>(args)...);
 }
 
-template <typename OPERATOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
-__global__ void call_operator_using_list(
-    DESCRIPTOR descriptor
-  , std::size_t count
-  , ARGS... args
-) {
-  const std::size_t index = blockIdx.x * blockDim.x + threadIdx.x;
-  if (!(index < count)) {
-      return;
-  }
-  OPERATOR::template apply<T,S>(descriptor, index, count, std::forward<ARGS>(args)...);
-}
-
 template <typename FUNCTOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
 __global__ void call_spatial_functor(
     LatticeView<DESCRIPTOR,S> lattice