diff --git a/test/integ/axpy/src/axpy.cpp b/test/integ/axpy/src/axpy.cpp
index 9b6c527f776a..605d6736b3ac 100644
--- a/test/integ/axpy/src/axpy.cpp
+++ b/test/integ/axpy/src/axpy.cpp
@@ -73,13 +73,15 @@ class AxpyKernel
     }
 };
 
-struct TestTemplate
-{
-template< typename TAcc >
-void operator()()
+using TestAccs = alpaka::test::acc::EnabledAccs<
+    alpaka::dim::DimInt<1u>,
+    std::size_t>;
+
+TEMPLATE_LIST_TEST_CASE( "axpy", "[axpy]", TestAccs)
 {
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
 #ifdef ALPAKA_CI
     Idx const numElements = 1u<<9u;
@@ -88,7 +90,7 @@ void operator()()
 #endif
 
     using Val = float;
-    using DevAcc = alpaka::dev::Dev<TAcc>;
+    using DevAcc = alpaka::dev::Dev<Acc>;
     using PltfAcc = alpaka::pltf::Pltf<DevAcc>;
     using QueueAcc = alpaka::test::queue::DefaultQueue<DevAcc>;
     using PltfHost = alpaka::pltf::PltfCpu;
@@ -112,7 +114,7 @@ void operator()()
 
     // Let alpaka calculate good block and grid sizes given our full problem extent.
     alpaka::workdiv::WorkDivMembers<Dim, Idx> const workDiv(
-        alpaka::workdiv::getValidWorkDiv<TAcc>(
+        alpaka::workdiv::getValidWorkDiv<Acc>(
             devAcc,
             extent,
             static_cast<Idx>(3u),
@@ -122,7 +124,7 @@ void operator()()
     std::cout
         << "AxpyKernel("
         << " numElements:" << numElements
-        << ", accelerator: " << alpaka::acc::getAccName<TAcc>()
+        << ", accelerator: " << alpaka::acc::getAccName<Acc>()
         << ", kernel: " << typeid(kernel).name()
         << ", workDiv: " << workDiv
         << ")" << std::endl;
@@ -181,7 +183,7 @@ void operator()()
 #endif
 
     // Create the kernel execution task.
-    auto const taskKernel(alpaka::kernel::createTaskKernel<TAcc>(
+    auto const taskKernel(alpaka::kernel::createTaskKernel<Acc>(
         workDiv,
         kernel,
         numElements,
@@ -218,13 +220,3 @@ void operator()()
 
     REQUIRE(resultCorrect);
 }
-};
-
-TEST_CASE( "axpy", "[axpy]")
-{
-    using TestAccs = alpaka::test::acc::EnabledAccs<
-        alpaka::dim::DimInt<1u>,
-        std::size_t>;
-
-    alpaka::meta::forEachType< TestAccs >( TestTemplate() );
-}
diff --git a/test/integ/mandelbrot/src/mandelbrot.cpp b/test/integ/mandelbrot/src/mandelbrot.cpp
index 8278a4acc88c..4031357289e7 100644
--- a/test/integ/mandelbrot/src/mandelbrot.cpp
+++ b/test/integ/mandelbrot/src/mandelbrot.cpp
@@ -302,13 +302,15 @@ auto writeTgaColorImage(
     }
 }
 
-struct TestTemplate
-{
-template< typename TAcc >
-void operator()()
+using TestAccs = alpaka::test::acc::EnabledAccs<
+    alpaka::dim::DimInt<2u>,
+    std::uint32_t>;
+
+TEMPLATE_LIST_TEST_CASE( "mandelbrot", "[mandelbrot]", TestAccs)
 {
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
 #ifdef ALPAKA_CI
     Idx const imageSize(1u<<5u);
@@ -324,7 +326,7 @@ void operator()()
     Idx const maxIterations(300u);
 
     using Val = std::uint32_t;
-    using DevAcc = alpaka::dev::Dev<TAcc>;
+    using DevAcc = alpaka::dev::Dev<Acc>;
     using PltfAcc = alpaka::pltf::Pltf<DevAcc>;
     using QueueAcc = alpaka::test::queue::DefaultQueue<DevAcc>;
     using PltfHost = alpaka::pltf::PltfCpu;
@@ -350,7 +352,7 @@ void operator()()
 
     // Let alpaka calculate good block and grid sizes given our full problem extent.
     alpaka::workdiv::WorkDivMembers<Dim, Idx> const workDiv(
-        alpaka::workdiv::getValidWorkDiv<TAcc>(
+        alpaka::workdiv::getValidWorkDiv<Acc>(
             devAcc,
             extent,
             alpaka::vec::Vec<Dim, Idx>::ones(),
@@ -362,7 +364,7 @@ void operator()()
         << " numRows:" << numRows
         << ", numCols:" << numCols
         << ", maxIterations:" << maxIterations
-        << ", accelerator: " << alpaka::acc::getAccName<TAcc>()
+        << ", accelerator: " << alpaka::acc::getAccName<Acc>()
         << ", kernel: " << typeid(kernel).name()
         << ", workDiv: " << workDiv
         << ")" << std::endl;
@@ -379,7 +381,7 @@ void operator()()
     alpaka::mem::view::copy(queue, bufColorAcc, bufColorHost, extent);
 
     // Create the kernel execution task.
-    auto const taskKernel(alpaka::kernel::createTaskKernel<TAcc>(
+    auto const taskKernel(alpaka::kernel::createTaskKernel<Acc>(
         workDiv,
         kernel,
         alpaka::mem::view::getPtrNative(bufColorAcc),
@@ -407,20 +409,10 @@ void operator()()
     alpaka::wait::wait(queue);
 
     // Write the image to a file.
-    std::string fileName("mandelbrot"+std::to_string(numCols)+"x"+std::to_string(numRows)+"_"+alpaka::acc::getAccName<TAcc>()+".tga");
+    std::string fileName("mandelbrot"+std::to_string(numCols)+"x"+std::to_string(numRows)+"_"+alpaka::acc::getAccName<Acc>()+".tga");
     std::replace(fileName.begin(), fileName.end(), '<', '_');
     std::replace(fileName.begin(), fileName.end(), '>', '_');
     writeTgaColorImage(
         fileName,
         bufColorHost);
 }
-};
-
-TEST_CASE( "mandelbrot", "[mandelbrot]")
-{
-    using TestAccs = alpaka::test::acc::EnabledAccs<
-        alpaka::dim::DimInt<2u>,
-        std::uint32_t>;
-
-    alpaka::meta::forEachType< TestAccs >( TestTemplate() );
-}
diff --git a/test/integ/matMul/src/matMul.cpp b/test/integ/matMul/src/matMul.cpp
index ca274bb9df30..cda1f0855ae3 100644
--- a/test/integ/matMul/src/matMul.cpp
+++ b/test/integ/matMul/src/matMul.cpp
@@ -199,13 +199,15 @@ namespace alpaka
     }
 }
 
-struct TestTemplate
-{
-template< typename TAcc >
-void operator()()
+using TestAccs = alpaka::test::acc::EnabledAccs<
+    alpaka::dim::DimInt<2u>,
+    std::uint32_t>;
+
+TEMPLATE_LIST_TEST_CASE( "matMul", "[matMul]", TestAccs)
 {
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
     Idx const m(64u);
     Idx const n(79u);
@@ -213,9 +215,9 @@ void operator()()
 
     using Val = std::uint32_t;
     using Vec2 = alpaka::vec::Vec<Dim, Idx>;
-    using DevAcc = alpaka::dev::Dev<TAcc>;
+    using DevAcc = alpaka::dev::Dev<Acc>;
     using PltfAcc = alpaka::pltf::Pltf<DevAcc>;
-    using QueueAcc = alpaka::test::queue::DefaultQueue<alpaka::dev::Dev<TAcc>>;
+    using QueueAcc = alpaka::test::queue::DefaultQueue<alpaka::dev::Dev<Acc>>;
     using PltfHost = alpaka::pltf::PltfCpu;
     using DevHost = alpaka::dev::Dev<PltfHost>;
     using QueueHost = alpaka::queue::QueueCpuNonBlocking;
@@ -255,7 +257,7 @@ void operator()()
 
     // Let alpaka calculate good block and grid sizes given our full problem extent.
     alpaka::workdiv::WorkDivMembers<Dim, Idx> const workDiv(
-        alpaka::workdiv::getValidWorkDiv<TAcc>(
+        alpaka::workdiv::getValidWorkDiv<Acc>(
             devAcc,
             extentC,
             alpaka::vec::Vec<Dim, Idx>::ones(),
@@ -267,7 +269,7 @@ void operator()()
         << "m:" << m
         << ", n:" << n
         << ", k:" << k
-        << ", accelerator: " << alpaka::acc::getAccName<TAcc>()
+        << ", accelerator: " << alpaka::acc::getAccName<Acc>()
         << ", kernel: " << typeid(kernel).name()
         << ", workDiv: " << workDiv
         << ")" << std::endl;
@@ -304,7 +306,7 @@ void operator()()
     alpaka::mem::view::copy(queueAcc, bufCAcc, bufCHost, extentC);
 
     // Create the kernel execution task.
-    auto const taskKernel(alpaka::kernel::createTaskKernel<TAcc>(
+    auto const taskKernel(alpaka::kernel::createTaskKernel<Acc>(
         workDiv,
         kernel,
         m,
@@ -353,13 +355,3 @@ void operator()()
 
     REQUIRE(resultCorrect);
 }
-};
-
-TEST_CASE( "matMul", "[matMul]")
-{
-    using TestAccs = alpaka::test::acc::EnabledAccs<
-        alpaka::dim::DimInt<2u>,
-        std::uint32_t>;
-
-    alpaka::meta::forEachType< TestAccs >( TestTemplate() );
-}
diff --git a/test/integ/separableCompilation/src/main.cpp b/test/integ/separableCompilation/src/main.cpp
index 1a0a1de138f5..317f4694b167 100644
--- a/test/integ/separableCompilation/src/main.cpp
+++ b/test/integ/separableCompilation/src/main.cpp
@@ -69,18 +69,20 @@ class SqrtKernel
     }
 };
 
-struct TestTemplate
-{
-template< typename TAcc >
-void operator()()
+using TestAccs = alpaka::test::acc::EnabledAccs<
+    alpaka::dim::DimInt<1u>,
+    std::size_t>;
+
+TEMPLATE_LIST_TEST_CASE( "separableCompilation", "[separableCompilation]", TestAccs)
 {
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Idx = alpaka::idx::Idx<Acc>;
 
     using Val = double;
 
-    using DevAcc = alpaka::dev::Dev<TAcc>;
+    using DevAcc = alpaka::dev::Dev<Acc>;
     using PltfAcc = alpaka::pltf::Pltf<DevAcc>;
-    using QueueAcc = alpaka::test::queue::DefaultQueue<alpaka::dev::Dev<TAcc>>;
+    using QueueAcc = alpaka::test::queue::DefaultQueue<alpaka::dev::Dev<Acc>>;
     using PltfHost = alpaka::pltf::PltfCpu;
     using DevHost = alpaka::dev::Dev<PltfHost>;
 
@@ -106,7 +108,7 @@ void operator()()
 
     // Let alpaka calculate good block and grid sizes given our full problem extent.
     alpaka::workdiv::WorkDivMembers<alpaka::dim::DimInt<1u>, Idx> const workDiv(
-        alpaka::workdiv::getValidWorkDiv<TAcc>(
+        alpaka::workdiv::getValidWorkDiv<Acc>(
             devAcc,
             extent,
             static_cast<Idx>(3u),
@@ -115,7 +117,7 @@ void operator()()
 
     std::cout
         << typeid(kernel).name() << "("
-        << "accelerator: " << alpaka::acc::getAccName<TAcc>()
+        << "accelerator: " << alpaka::acc::getAccName<Acc>()
         << ", workDiv: " << workDiv
         << ", numElements:" << numElements
         << ")" << std::endl;
@@ -142,7 +144,7 @@ void operator()()
     alpaka::mem::view::copy(queueAcc, memBufAccB, memBufHostB, extent);
 
     // Create the executor task.
-    auto const taskKernel(alpaka::kernel::createTaskKernel<TAcc>(
+    auto const taskKernel(alpaka::kernel::createTaskKernel<Acc>(
         workDiv,
         kernel,
         alpaka::mem::view::getPtrNative(memBufAccA),
@@ -180,13 +182,3 @@ void operator()()
 
     REQUIRE(true == resultCorrect);
 }
-};
-
-TEST_CASE( "separableCompilation", "[separableCompilation]")
-{
-    using TestAccs = alpaka::test::acc::EnabledAccs<
-        alpaka::dim::DimInt<1u>,
-        std::size_t>;
-
-    alpaka::meta::forEachType< TestAccs >( TestTemplate() );
-}
diff --git a/test/integ/sharedMem/src/sharedMem.cpp b/test/integ/sharedMem/src/sharedMem.cpp
index 842084764aef..e1e49b59a108 100644
--- a/test/integ/sharedMem/src/sharedMem.cpp
+++ b/test/integ/sharedMem/src/sharedMem.cpp
@@ -134,20 +134,22 @@ namespace alpaka
     }
 }
 
-struct TestTemplate
-{
-template< typename TAcc >
-void operator()()
+using TestAccs = alpaka::test::acc::EnabledAccs<
+    alpaka::dim::DimInt<1u>,
+    std::uint32_t>;
+
+TEMPLATE_LIST_TEST_CASE( "sharedMem", "[sharedMem]", TestAccs)
 {
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
     Idx const numElements = 1u<<16u;
 
     using Val = std::int32_t;
     using TnumUselessWork = std::integral_constant<Idx, 100>;
 
-    using DevAcc = alpaka::dev::Dev<TAcc>;
+    using DevAcc = alpaka::dev::Dev<Acc>;
     using PltfAcc = alpaka::pltf::Pltf<DevAcc>;
     using QueueAcc = alpaka::test::queue::DefaultQueue<DevAcc>;
 
@@ -165,7 +167,7 @@ void operator()()
 
     // Set the grid blocks extent.
     alpaka::workdiv::WorkDivMembers<Dim, Idx> const workDiv(
-        alpaka::workdiv::getValidWorkDiv<TAcc>(
+        alpaka::workdiv::getValidWorkDiv<Acc>(
             devAcc,
             numElements,
             static_cast<Idx>(1u),
@@ -174,7 +176,7 @@ void operator()()
 
     std::cout
         << "SharedMemKernel("
-        << " accelerator: " << alpaka::acc::getAccName<TAcc>()
+        << " accelerator: " << alpaka::acc::getAccName<Acc>()
         << ", kernel: " << typeid(kernel).name()
         << ", workDiv: " << workDiv
         << ")" << std::endl;
@@ -193,7 +195,7 @@ void operator()()
     alpaka::mem::view::copy(queue, blockRetValsAcc, blockRetVals, resultElemCount);
 
     // Create the kernel execution task.
-    auto const taskKernel(alpaka::kernel::createTaskKernel<TAcc>(
+    auto const taskKernel(alpaka::kernel::createTaskKernel<Acc>(
         workDiv,
         kernel,
         alpaka::mem::view::getPtrNative(blockRetValsAcc)));
@@ -229,13 +231,3 @@ void operator()()
 
     REQUIRE(resultCorrect);
 }
-};
-
-TEST_CASE( "sharedMem", "[sharedMem]")
-{
-    using TestAccs = alpaka::test::acc::EnabledAccs<
-        alpaka::dim::DimInt<1u>,
-        std::uint32_t>;
-
-    alpaka::meta::forEachType< TestAccs >( TestTemplate() );
-}
diff --git a/test/unit/kernel/src/KernelStdFunction.cpp b/test/unit/kernel/src/KernelStdFunction.cpp
index 234858ac5883..0856ea3cd81b 100644
--- a/test/unit/kernel/src/KernelStdFunction.cpp
+++ b/test/unit/kernel/src/KernelStdFunction.cpp
@@ -22,9 +22,9 @@
 
 //-----------------------------------------------------------------------------
 template<
-    typename TAcc>
+    typename Acc>
 void ALPAKA_FN_ACC kernelFn(
-    TAcc const & acc,
+    Acc const & acc,
     bool * success,
     std::int32_t val)
 {
@@ -36,47 +36,31 @@ void ALPAKA_FN_ACC kernelFn(
 // std::function and std::bind is only allowed on CPU
 #if !BOOST_LANG_CUDA && !BOOST_LANG_HIP
 //-----------------------------------------------------------------------------
-struct TestTemplateStdFunction
+TEMPLATE_LIST_TEST_CASE( "stdFunctionKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
 {
-template< typename TAcc >
-void operator()()
-{
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-    alpaka::test::KernelExecutionFixture<TAcc> fixture(
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
         alpaka::vec::Vec<Dim, Idx>::ones());
 
-    const auto kernel = std::function<void(TAcc const &, bool *, std::int32_t)>( kernelFn<TAcc> );
+    const auto kernel = std::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
     REQUIRE(fixture(kernel, 42));
-  }
-};
-
-TEST_CASE( "stdFunctionKernelIsWorking", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateStdFunction() );
 }
 
 //-----------------------------------------------------------------------------
-struct TestTemplateStdBind
+TEMPLATE_LIST_TEST_CASE( "stdBindKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
 {
-template< typename TAcc >
-void operator()()
-{
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-    alpaka::test::KernelExecutionFixture<TAcc> fixture(
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
         alpaka::vec::Vec<Dim, Idx>::ones());
 
-    const auto kernel = std::bind( kernelFn<TAcc>, std::placeholders::_1, std::placeholders::_2, 42 );
+    const auto kernel = std::bind( kernelFn<Acc>, std::placeholders::_1, std::placeholders::_2, 42 );
     REQUIRE(fixture(kernel));
-  }
-};
-
-TEST_CASE( "stdBindKernelIsWorking", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateStdBind() );
 }
 #endif
 
@@ -84,28 +68,21 @@ TEST_CASE( "stdBindKernelIsWorking", "[kernel]")
 #if 0
 //#if BOOST_LANG_CUDA
 // clang as a native CUDA compiler does not seem to support nvstd::function when ALPAKA_ACC_GPU_CUDA_ONLY_MODE is used.
-// error: reference to __device__ function 'kernelFn<alpaka::acc::AccGpuCudaRt<std::__1::integral_constant<unsigned long, 1>, unsigned long> >' in __host__ function const auto kernel = nvstd::function<void(TAcc const &, bool *, std::int32_t)>( kernelFn<TAcc> );
+// error: reference to __device__ function 'kernelFn<alpaka::acc::AccGpuCudaRt<std::__1::integral_constant<unsigned long, 1>, unsigned long> >' in __host__ function const auto kernel = nvstd::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
 #if !(defined(ALPAKA_ACC_GPU_CUDA_ONLY_MODE) && BOOST_COMP_CLANG_CUDA)
 //-----------------------------------------------------------------------------
-struct TestTemplateNvstdFunction
-{
-template< typename TAcc >
-void operator()()
+TEMPLATE_LIST_TEST_CASE( "nvstdFunctionKernelIsWorking", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    using Dim = alpaka::dim::Dim<TAcc>;
-    using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-    alpaka::test::KernelExecutionFixture<TAcc> fixture(
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
         alpaka::vec::Vec<Dim, Idx>::ones());
 
-    const auto kernel = nvstd::function<void(TAcc const &, bool *, std::int32_t)>( kernelFn<TAcc> );
+    const auto kernel = nvstd::function<void(Acc const &, bool *, std::int32_t)>( kernelFn<Acc> );
     REQUIRE(fixture(kernel, 42));
-  }
-};
-
-TEST_CASE( "nvstdFunctionKernelIsWorking", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateNvstdFunction() );
 }
+
 #endif
 #endif
diff --git a/test/unit/kernel/src/KernelWithTemplateArgumentDeduction.cpp b/test/unit/kernel/src/KernelWithTemplateArgumentDeduction.cpp
index 7695ee53c71d..4873e5763db2 100644
--- a/test/unit/kernel/src/KernelWithTemplateArgumentDeduction.cpp
+++ b/test/unit/kernel/src/KernelWithTemplateArgumentDeduction.cpp
@@ -25,12 +25,12 @@ class KernelInvocationTemplateDeductionValueSemantics
     //-----------------------------------------------------------------------------
     ALPAKA_NO_HOST_ACC_WARNING
         template<
-        typename TAcc,
+        typename Acc,
         typename TByValue,
         typename TByConstValue,
         typename TByConstReference>
         ALPAKA_FN_ACC auto operator()(
-            TAcc const & acc,
+            Acc const & acc,
             bool * success,
             TByValue,
             TByConstValue const,
@@ -39,7 +39,7 @@ class KernelInvocationTemplateDeductionValueSemantics
     {
         ALPAKA_CHECK(
             *success,
-            static_cast<alpaka::idx::Idx<TAcc>>(1) == (alpaka::workdiv::getWorkDiv<alpaka::Grid, alpaka::Threads>(acc)).prod());
+            static_cast<alpaka::idx::Idx<Acc>>(1) == (alpaka::workdiv::getWorkDiv<alpaka::Grid, alpaka::Threads>(acc)).prod());
 
         static_assert(
             std::is_same<TByValue, TExpected>::value,
@@ -55,63 +55,54 @@ class KernelInvocationTemplateDeductionValueSemantics
 };
 
 //-----------------------------------------------------------------------------
-struct TestTemplateDeductionFromValue
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromValue", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionValueSemantics< Value > kernel;
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionValueSemantics< Value > kernel;
 
-        Value value{ };
-        REQUIRE(fixture(kernel, value, value, value));
-    }
-};
+    Value value{ };
+    REQUIRE(fixture(kernel, value, value, value));
+}
 
-struct TestTemplateDeductionFromConstValue
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromConstValue", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionValueSemantics< Value > kernel;
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionValueSemantics< Value > kernel;
 
-        Value const constValue{ };
-        REQUIRE(fixture(kernel, constValue, constValue, constValue));
-    }
-};
+    Value const constValue{ };
+    REQUIRE(fixture(kernel, constValue, constValue, constValue));
+}
 
-struct TestTemplateDeductionFromConstReference
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromConstReference", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionValueSemantics< Value > kernel;
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionValueSemantics< Value > kernel;
 
-        Value value{ };
-        Value const & constReference = value;
-        REQUIRE(fixture(kernel, constReference, constReference, constReference));
-    }
-};
+    Value value{ };
+    Value const & constReference = value;
+    REQUIRE(fixture(kernel, constReference, constReference, constReference));
+}
 
 //#############################################################################
 template<
@@ -124,11 +115,11 @@ class KernelInvocationTemplateDeductionPointerSemantics
     //-----------------------------------------------------------------------------
     ALPAKA_NO_HOST_ACC_WARNING
         template<
-        typename TAcc,
+        typename Acc,
         typename TByPointer,
         typename TByPointerToConst>
         ALPAKA_FN_ACC auto operator()(
-            TAcc const & acc,
+            Acc const & acc,
             bool * success,
             TByPointer *,
             TByPointerToConst const *) const
@@ -136,7 +127,7 @@ class KernelInvocationTemplateDeductionPointerSemantics
     {
         ALPAKA_CHECK(
             *success,
-            static_cast<alpaka::idx::Idx<TAcc>>(1) == (alpaka::workdiv::getWorkDiv<alpaka::Grid, alpaka::Threads>(acc)).prod());
+            static_cast<alpaka::idx::Idx<Acc>>(1) == (alpaka::workdiv::getWorkDiv<alpaka::Grid, alpaka::Threads>(acc)).prod());
 
         static_assert(
             std::is_same<TByPointer, TExpectedFirst>::value,
@@ -149,115 +140,68 @@ class KernelInvocationTemplateDeductionPointerSemantics
 };
 
 //-----------------------------------------------------------------------------
-struct TestTemplateDeductionFromPointer
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromPointer", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
-
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionPointerSemantics< Value > kernel;
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-        Value value{ };
-        Value * pointer = &value;
-        REQUIRE(fixture(kernel, pointer, pointer));
-    }
-};
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionPointerSemantics< Value > kernel;
 
-struct TestTemplateDeductionFromPointerToConst
-{
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
-
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
-
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionPointerSemantics< Value const, Value > kernel;
-
-        Value const constValue{ };
-        Value const * pointerToConst = &constValue;
-        REQUIRE(fixture(kernel, pointerToConst, pointerToConst));
-    }
-};
+    Value value{ };
+    Value * pointer = &value;
+    REQUIRE(fixture(kernel, pointer, pointer));
+}
 
-struct TestTemplateDeductionFromStaticArray
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromPointerToConst", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionPointerSemantics< Value > kernel;
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionPointerSemantics< Value const, Value > kernel;
 
-        Value staticArray[4] = { };
-        REQUIRE(fixture(kernel, staticArray, staticArray));
-    }
-};
+    Value const constValue{ };
+    Value const * pointerToConst = &constValue;
+    REQUIRE(fixture(kernel, pointerToConst, pointerToConst));
+}
 
-struct TestTemplateDeductionFromConstStaticArray
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromStaticArray", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    template< typename TAcc >
-    void operator()()
-    {
-        using Dim = alpaka::dim::Dim<TAcc>;
-        using Idx = alpaka::idx::Idx<TAcc>;
-
-        alpaka::test::KernelExecutionFixture<TAcc> fixture(
-            alpaka::vec::Vec<Dim, Idx>::ones());
-
-        using Value = std::int32_t;
-        KernelInvocationTemplateDeductionPointerSemantics< Value const, Value > kernel;
-
-        Value const constStaticArray[4] = { };
-        REQUIRE(fixture(kernel, constStaticArray, constStaticArray));
-    }
-};
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromValue", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromValue() );
-}
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromConstValue", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromConstValue() );
-}
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionPointerSemantics< Value > kernel;
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromConstReference", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromConstReference() );
+    Value staticArray[4] = { };
+    REQUIRE(fixture(kernel, staticArray, staticArray));
 }
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromPointer", "[kernel]")
+TEMPLATE_LIST_TEST_CASE( "kernelFuntionObjectTemplateDeductionFromConstStaticArray", "[kernel]", alpaka::test::acc::TestAccs)
 {
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromPointer() );
-}
+    using Acc = TestType;
+    using Dim = alpaka::dim::Dim<Acc>;
+    using Idx = alpaka::idx::Idx<Acc>;
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromPointerToConst", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromPointerToConst() );
-}
+    alpaka::test::KernelExecutionFixture<Acc> fixture(
+        alpaka::vec::Vec<Dim, Idx>::ones());
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromStaticArray", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromStaticArray() );
-}
+    using Value = std::int32_t;
+    KernelInvocationTemplateDeductionPointerSemantics< Value const, Value > kernel;
 
-TEST_CASE( "kernelFuntionObjectTemplateDeductionFromConstStaticArray", "[kernel]")
-{
-    alpaka::meta::forEachType< alpaka::test::acc::TestAccs >( TestTemplateDeductionFromConstStaticArray() );
+    Value const constStaticArray[4] = { };
+    REQUIRE(fixture(kernel, constStaticArray, constStaticArray));
 }