summaryrefslogtreecommitdiff
path: root/tangle/LLBM/kernel
diff options
context:
space:
mode:
authorAdrian Kummerlaender2021-09-12 14:01:55 +0200
committerAdrian Kummerlaender2021-09-12 14:01:55 +0200
commit32dd41a728ce10113032e20955ba08f8de449857 (patch)
tree8a5f851a342590d05137b2c900a7d3c4e8c545b6 /tangle/LLBM/kernel
parent8bca21a550e0ef134d51c4c4c007720885d76791 (diff)
downloadLiterateLB-32dd41a728ce10113032e20955ba08f8de449857.tar
LiterateLB-32dd41a728ce10113032e20955ba08f8de449857.tar.gz
LiterateLB-32dd41a728ce10113032e20955ba08f8de449857.tar.bz2
LiterateLB-32dd41a728ce10113032e20955ba08f8de449857.tar.lz
LiterateLB-32dd41a728ce10113032e20955ba08f8de449857.tar.xz
LiterateLB-32dd41a728ce10113032e20955ba08f8de449857.tar.zst
LiterateLB-32dd41a728ce10113032e20955ba08f8de449857.zip
Start using C++ cuda-api-wrapper instead of raw CUDA
Diffstat (limited to 'tangle/LLBM/kernel')
-rw-r--r--tangle/LLBM/kernel/executor.h37
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