From 32dd41a728ce10113032e20955ba08f8de449857 Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Sun, 12 Sep 2021 14:01:55 +0200 Subject: Start using C++ cuda-api-wrapper instead of raw CUDA --- tangle/LLBM/kernel/executor.h | 37 ------------------------------------- 1 file changed, 37 deletions(-) (limited to 'tangle/LLBM/kernel') 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 -__global__ void call_operator( - LatticeView 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(DESCRIPTOR(), f_curr, f_next, gid, std::forward(args)...); - for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) { - *preshifted_f[iPop] = f_next[iPop]; - } -} - template __global__ void call_functor( LatticeView lattice @@ -110,19 +86,6 @@ __global__ void call_operator_using_list( OPERATOR::template apply(lattice, index, count, std::forward(args)...); } -template -__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(descriptor, index, count, std::forward(args)...); -} - template __global__ void call_spatial_functor( LatticeView lattice -- cgit v1.2.3