SYCL: Shorten mangled kernel name types

[alexxy/gromacs.git] / src / gromacs / mdlib / gpuforcereduction_impl_internal_sycl.cpp

diff --git a/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp b/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp
index 3bc21c43de8ccfd860890d8f383aad18d0ec0397..e2ee7b46979d1bd12dccf8600faeed8f19553e51 100644 (file)
--- a/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp
+++ b/src/gromacs/mdlib/gpuforcereduction_impl_internal_sycl.cpp
@@ -54,6 +54,10 @@
 #include "gromacs/gpu_utils/gpueventsynchronizer_sycl.h"
 #include "gromacs/utility/template_mp.h"
 
+//! \brief Class name for reduction kernel
+template<bool addRvecForce, bool accumulateForce>
+class ReduceKernel;
+
 namespace gmx
 {
 
@@ -93,9 +97,6 @@ static auto reduceKernel(cl::sycl::handler&                                 cgh,
     };
 }
 
-template<bool addRvecForce, bool accumulateForce>
-class ReduceKernelName;
-
 template<bool addRvecForce, bool accumulateForce>
 static void launchReductionKernel_(const int                   numAtoms,
                                    const int                   atomStart,
@@ -114,7 +115,7 @@ static void launchReductionKernel_(const int                   numAtoms,
     queue.submit([&](cl::sycl::handler& cgh) {
         auto kernel = reduceKernel<addRvecForce, accumulateForce>(
                 cgh, b_nbnxmForce, b_rvecForceToAdd, b_forceTotal, b_cell, atomStart);
-        cgh.parallel_for<ReduceKernelName<addRvecForce, accumulateForce>>(rangeNumAtoms, kernel);
+        cgh.parallel_for<ReduceKernel<addRvecForce, accumulateForce>>(rangeNumAtoms, kernel);
     });
 }