small adjustments and flatter init optimization

Author: psychocoderHPC

examples/getAvailableSlots/source/main.cpp

@@ -123,19 +123,22 @@ auto main() -> int
             if(result != EXIT_SUCCESS)
                 return;
 
+            std::cout << alpaka::onHost::demangledName<FlatterScatter<FlatterScatterHeapConfig>>() << ":\n";
             result = runExample<
                 Executor,
                 FlatterScatter<FlatterScatterHeapConfig>,
                 mallocMC::ReservePoolPolicies::AlpakaBuf>(deviceSpec, exec);
             if(result != EXIT_SUCCESS)
                 return;
+            std::cout << alpaka::onHost::demangledName<Scatter<FlatterScatterHeapConfig>>() << ":\n";
             result = runExample<Executor, Scatter<FlatterScatterHeapConfig>, mallocMC::ReservePoolPolicies::AlpakaBuf>(
                 deviceSpec,
                 exec);
 #if ALPAKA_LANG_CUDA
 #    ifdef mallocMC_HAS_Gallatin_AVAILABLE
             if(result == EXIT_SUCCESS)
             {
+                std::cout << alpaka::onHost::demangledName<mallocMC::CreationPolicies::GallatinCuda<>>() << ":\n";
                 result = runExample<
                     Executor,
                     mallocMC::CreationPolicies::GallatinCuda<>,
@@ -145,7 +148,10 @@ auto main() -> int
 #    endif
 #endif
             if(result == EXIT_SUCCESS)
+            {
+                std::cout << alpaka::onHost::demangledName<OldMalloc>() << ":\n";
                 result = runExample<Executor, OldMalloc, mallocMC::ReservePoolPolicies::Noop>(deviceSpec, exec);
+            }
         },
         alpaka::onHost::allBackends(alpaka::onHost::enabledDeviceSpecs, alpaka::exec::enabledExecutors));
     return result;

examples/native-cuda/source/main.cu

@@ -26,11 +26,13 @@
   THE SOFTWARE.
 */
 
+#include "mallocMC/span.hpp"
+
 #include <mallocMC/mallocMC.cuh>
 
 #include <cstdint>
 #include <cstdlib>
-#include <iostream>
+#include <functional>
 
 /**
  * @brief Computes the sum of squares of the first `n` natural numbers.
@@ -65,26 +67,19 @@ __device__ auto sumOfSquares(auto const n)
  */
 __global__ void oneDotProductPerThread(mallocMC::CudaMemoryManager<> memoryManager, uint64_t numValues)
 {
+    using mallocMC::span;
     uint64_t tid = threadIdx.x + blockIdx.x * blockDim.x;
 
     // Not very realistic, all threads are doing this on their own:
-    auto* a = reinterpret_cast<uint64_t*>(memoryManager.malloc(numValues * sizeof(uint64_t)));
-    auto* b = reinterpret_cast<uint64_t*>(memoryManager.malloc(numValues * sizeof(uint64_t)));
-    if(a == nullptr || b == nullptr)
-    {
-        printf("Thread %lu: device allocation failed.\n", tid);
-        __trap();
-    }
+    auto a
+        = span<uint64_t>(reinterpret_cast<uint64_t*>(memoryManager.malloc(numValues * sizeof(uint64_t))), numValues);
+    auto b
+        = span<uint64_t>(reinterpret_cast<uint64_t*>(memoryManager.malloc(numValues * sizeof(uint64_t))), numValues);
 
-    for(uint64_t i = 0; i < numValues; ++i)
-    {
-        a[i] = tid + i;
-        b[i] = tid + i;
-    }
+    std::iota(std::begin(a), std::end(a), tid);
+    std::iota(std::begin(b), std::end(b), tid);
 
-    uint64_t result = 0U;
-    for(uint64_t i = 0; i < numValues; ++i)
-        result += a[i] * b[i];
+    uint64_t result = std::transform_reduce(std::cbegin(a), std::cend(a), std::cbegin(b), 0U);
 
     auto expected = sumOfSquares(numValues + tid - 1) - (tid > 0 ? sumOfSquares(tid - 1) : 0);
     if(result != expected)
@@ -93,8 +88,8 @@ __global__ void oneDotProductPerThread(mallocMC::CudaMemoryManager<> memoryManag
         __trap();
     }
 
-    memoryManager.free(a);
-    memoryManager.free(b);
+    memoryManager.free(a.data());
+    memoryManager.free(b.data());
 }
 
 int main()
@@ -106,5 +101,4 @@ int main()
 
     std::cout << "Running native CUDA kernel." << std::endl;
     oneDotProductPerThread<<<8, 256>>>(memoryManager, numValues);
-    cudaDeviceSynchronize();
 }

examples/vectorAdd/source/main.cpp

@@ -49,8 +49,12 @@ struct ShrinkConfig
 struct VectorAddKernel
 {
     template<typename TAcc, typename TAllocHandle, typename TSums>
-    ALPAKA_FN_ACC void operator()(TAcc const& acc, TAllocHandle allocHandle, TSums sums, std::uint32_t len, std::uint32_t count)
-        const
+    ALPAKA_FN_ACC void operator()(
+        TAcc const& acc,
+        TAllocHandle allocHandle,
+        TSums sums,
+        std::uint32_t len,
+        std::uint32_t count) const
     {
         for(auto [id] : alpaka::onAcc::makeIdxMap(acc, alpaka::onAcc::worker::threadsInGrid, alpaka::IdxRange{count}))
         {
@@ -121,10 +125,8 @@ auto runExample(auto const& deviceSpec, TExecutor exec) -> int
     std::cout << Allocator::info("\n") << '\n';
 
     auto frameExtent = alpaka::Vec{Idx{threadsPerBlock}};
-    auto frameSpec = alpaka::onHost::FrameSpec{
-        alpaka::divCeil(alpaka::Vec{Idx{numWorkers}}, frameExtent),
-        frameExtent,
-        exec};
+    auto frameSpec
+        = alpaka::onHost::FrameSpec{alpaka::divCeil(alpaka::Vec{Idx{numWorkers}}, frameExtent), frameExtent, exec};
     queue.enqueue(
         frameSpec,
         alpaka::KernelBundle{VectorAddKernel{}, alloc.getAllocatorHandle(), sumsAcc, localLength, numWorkers});

include/mallocMC/creationPolicies/FlatterScatter.hpp

@@ -81,18 +81,13 @@ namespace mallocMC::CreationPolicies::FlatterScatterAlloc
         ALPAKA_FN_INLINE ALPAKA_FN_ACC static auto init(auto const& acc, void* accessBlocksPointer, auto heapSize)
             -> void
         {
-            auto const threadsInGrid = acc.getExtentsOf(alpaka::onAcc::origin::grid, alpaka::onAcc::unit::threads);
-            auto const numThreads = threadsInGrid.product();
-            auto const idx = static_cast<uint32_t>(alpaka::linearize(
-                threadsInGrid,
-                acc.getIdxWithin(alpaka::onAcc::origin::grid, alpaka::onAcc::unit::threads)));
             auto* accessBlocks = static_cast<MyAccessBlock*>(accessBlocksPointer);
 
-            for(uint32_t i = idx; i < numBlocks(heapSize) * MyAccessBlock::numPages(); i += numThreads)
+            for(auto [blockIdx, pageIdx] : alpaka::onAcc::makeIdxMap(
+                    acc,
+                    alpaka::onAcc::worker::threadsInGrid,
+                    alpaka::IdxRange{alpaka::Vec{numBlocks(heapSize), MyAccessBlock::numPages()}}))
             {
-                auto blockIdx = i / MyAccessBlock::numPages();
-                auto pageIdx = i % MyAccessBlock::numPages();
-
                 accessBlocks[blockIdx].init(acc, pageIdx);
             }
         }
@@ -343,7 +338,6 @@ namespace mallocMC::CreationPolicies::FlatterScatterAlloc
             Heap<T_HeapConfig, T_HashConfig, T_AlignmentPolicy>::init(acc, m_heapmem, m_memsize);
         }
     };
-
 } // namespace mallocMC::CreationPolicies::FlatterScatterAlloc
 
 namespace mallocMC::CreationPolicies
@@ -421,9 +415,11 @@ namespace mallocMC::CreationPolicies
                 return;
             }
             auto numPagesPerBlock = MyHeap::MyAccessBlock::numPages();
-            queue.enqueue(
-                alpaka::onHost::FrameSpec{alpaka::Vec{numBlocks}, alpaka::Vec{numPagesPerBlock}, TExecutor{}},
-                alpaka::KernelBundle{FlatterScatterAlloc::InitKernel{}, heap, pool, memsize});
+            auto frameSpec = alpaka::onHost::getFrameSpec<uint32_t>(
+                queue.getDevice(),
+                TExecutor{},
+                alpaka::Vec{numBlocks, numPagesPerBlock});
+            queue.enqueue(frameSpec, alpaka::KernelBundle{FlatterScatterAlloc::InitKernel{}, heap, pool, memsize});
             alpaka::onHost::wait(queue);
         }
 
@@ -479,6 +475,4 @@ namespace mallocMC::CreationPolicies
             using HashConfig = T_HashConfig;
         };
     };
-
-
 } // namespace mallocMC::CreationPolicies