use PMACC_TYPEKERNEL

- use `PMACC_TYPEKERNEL` macro to start a kernel
- tranform global kernel function to a functor

Author: psychocoderHPC

src/libPMacc/test/particles/IdProvider.hpp

@@ -41,15 +41,20 @@ namespace bmpl = boost::mpl;
 
 namespace
 {
-    template<class T_IdProvider, class T_Box>
-    __global__ void generateIds(T_Box outputbox, uint32_t numThreads, uint32_t numIdsPerThread)
+    template< typename T_IdProvider >
+    struct generateIds
     {
-        const uint32_t localId = blockIdx.x * blockDim.x + threadIdx.x;
-        if(localId >= numThreads)
-            return;
-        for(uint32_t i=0; i<numIdsPerThread; i++)
-            outputbox(i * numThreads + localId) = T_IdProvider::getNewId();
-    }
+        template<class T_Box>
+        DINLINE void operator()(T_Box outputbox, uint32_t numThreads, uint32_t numIdsPerThread) const
+        {
+            const uint32_t localId = blockIdx.x * blockDim.x + threadIdx.x;
+            if(localId < numThreads)
+            {
+                for(uint32_t i=0; i<numIdsPerThread; i++)
+                    outputbox(i * numThreads + localId) = T_IdProvider::getNewId();
+            }
+        }
+    };
 }
 
 /**
@@ -115,7 +120,7 @@ struct IdProviderTest
         BOOST_REQUIRE_EQUAL(IdProvider::getNewId(), state.nextId);
         // Generate the same IDs on the device
         PMacc::HostDeviceBuffer<uint64_t, 1> idBuf(numIds);
-        __cudaKernel(generateIds<IdProvider>)(numBlocks, numThreadsPerBlock)
+        PMACC_TYPEKERNEL(generateIds<IdProvider>)(numBlocks, numThreadsPerBlock)
                 (idBuf.getDeviceBuffer().getDataBox(), numThreads, numIdsPerThread);
         idBuf.deviceToHost();
         BOOST_REQUIRE_EQUAL(numIds, ids.size());

src/libPMacc/test/random/2DDistribution.cu

@@ -38,17 +38,20 @@
 typedef PMacc::DataSpace<DIM2> Space2D;
 typedef PMacc::DataSpace<DIM3> Space3D;
 
-template<class T_DataBox, class T_Random>
-__global__ void RandomFiller(T_DataBox box, Space2D boxSize, T_Random rand, uint32_t numSamples)
+struct RandomFiller
 {
-    const Space3D ownIdx = Space3D(threadIdx) + Space3D(blockIdx) * Space3D(blockDim);
-    rand.init(ownIdx.shrink<2>());
-    for(uint32_t i=0; i<numSamples; i++)
+    template<class T_DataBox, class T_Random>
+    DINLINE void operator()(T_DataBox box, Space2D boxSize, T_Random rand, uint32_t numSamples) const
     {
-        Space2D idx = rand(boxSize);
-        atomicAdd(&box(idx), 1);
+        const Space3D ownIdx = Space3D(threadIdx) + Space3D(blockIdx) * Space3D(blockDim);
+        rand.init(ownIdx.shrink<2>());
+        for(uint32_t i=0; i<numSamples; i++)
+        {
+            Space2D idx = rand(boxSize);
+            atomicAdd(&box(idx), 1);
+        }
     }
-}
+};
 
 template<class T_RNGProvider>
 struct GetRandomIdx
@@ -130,7 +133,7 @@ void generateRandomNumbers(const Space2D& rngSize, uint32_t numSamples, T_Device
     Space2D gridSize(rngSize / blockSize);
 
     CUDA_CHECK(cudaEventRecord(start));
-    __cudaKernel(RandomFiller)(gridSize, blockSize)(buffer.getDataBox(), buffer.getDataSpace(), rand, numSamples);
+    PMACC_TYPEKERNEL(RandomFiller)(gridSize, blockSize)(buffer.getDataBox(), buffer.getDataSpace(), rand, numSamples);
     CUDA_CHECK(cudaEventRecord(stop));
     CUDA_CHECK(cudaEventSynchronize(stop));
     float milliseconds = 0;