* removed check for CUDA compute capability, since capability 3 is re…

…quired since CUDA 9 * removed cudaSetDeviceFlags, as it's not needed * added copyright and name in SimpleMalloc.hpp * removed 3 obsolete comments * improved declaration of devAllocatorStorage buffer in Allocator
alpaka-group · May 20, 2020 · 29de154 · 29de154
1 parent 651aa92
commit 29de154
Show file tree

Hide file tree

Showing 6 changed files with 14 additions and 53 deletions.
diff --git a/examples/mallocMC_example01.cpp b/examples/mallocMC_example01.cpp
@@ -82,7 +82,7 @@ ALPAKA_STATIC_ACC_MEM_GLOBAL int ** arA;
 ALPAKA_STATIC_ACC_MEM_GLOBAL int ** arB;
 ALPAKA_STATIC_ACC_MEM_GLOBAL int ** arC;
 
-void run()
+auto main() -> int
 {
     constexpr auto block = 32;
     constexpr auto grid = 32;
@@ -243,28 +243,6 @@ void run()
             alpaka::kernel::createTaskKernel<Acc>(
                 workDiv, freeArrayPointers, scatterAlloc.getAllocatorHandle()));
     }
-}
-
-auto main() -> int
-{
-#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
-    int computeCapabilityMajor = 0;
-    cudaDeviceGetAttribute(
-        &computeCapabilityMajor, cudaDevAttrComputeCapabilityMajor, 0);
-    int computeCapabilityMinor = 0;
-    cudaDeviceGetAttribute(
-        &computeCapabilityMinor, cudaDevAttrComputeCapabilityMinor, 0);
-
-    if(computeCapabilityMajor < 2)
-    {
-        std::cerr << "Error: Compute Capability >= 2.0 required. (is "
-                  << computeCapabilityMajor << "." << computeCapabilityMinor
-                  << ")\n";
-        return 1;
-    }
-#endif
-
-    run();
 
     return 0;
-}
+}
diff --git a/src/include/mallocMC/allocator.hpp b/src/include/mallocMC/allocator.hpp
@@ -140,12 +140,11 @@ namespace mallocMC
     private:
         std::function<void()>
             freeLastAlloc; // calles free(dev), with the dev used during alloc
-        decltype(alpaka::mem::buf::alloc<
-                 DevAllocator,
-                 alpaka::idx::traits::IdxType<AlpakaAcc>::type>(
-            std::declval<const alpaka::dev::Dev<AlpakaAcc>>(),
-            std::declval<
-                const alpaka::idx::traits::IdxType<AlpakaAcc>::type>()))
+        alpaka::mem::buf::Buf<
+            alpaka::dev::Dev<AlpakaAcc>,
+            DevAllocator,
+            typename alpaka::dim::traits::DimType<AlpakaAcc>::type,
+            typename alpaka::idx::traits::IdxType<AlpakaAcc>::type>
             devAllocatorStorage; // TODO(bgruber): since we want to have an
                                  // empty/invalid buffer sometimes, this should
                                  // be std::optional<>

diff --git a/src/include/mallocMC/creationPolicies/Scatter_impl.hpp b/src/include/mallocMC/creationPolicies/Scatter_impl.hpp
@@ -1278,7 +1278,6 @@ namespace mallocMC
 
                 // printf("Block %d, id %d: activeThreads=%d
                 // linearId=%d\n",blockIdx.x,threadIdx.x,activeThreads,linearId);
-                // // alpaka
                 const unsigned temp = getAvailaibleSlotsDeviceFunction(
                     acc, slotSize, linearId, activeThreads);
                 if(temp)

diff --git a/src/include/mallocMC/mallocMC_utils.hpp b/src/include/mallocMC/mallocMC_utils.hpp
@@ -74,7 +74,7 @@ namespace mallocMC
         return __lane_id();
 #elif defined(_CUDA_ARCH_)
         std::uint32_t mylaneid;
-        asm("mov.u32 %0, %%laneid;" : "=r"(mylaneid)); // alpaka
+        asm("mov.u32 %0, %%laneid;" : "=r"(mylaneid));
         return mylaneid;
 #else
         return 0u;
@@ -118,7 +118,7 @@ namespace mallocMC
         asm("mov.u32 %0, %%smid;" : "=r"(mysmid));
         return mysmid;
 #else
-        return 0u; // TODO(bgruber)
+        return 0u;
 #endif
     }
 

diff --git a/src/include/mallocMC/reservePoolPolicies/SimpleMalloc.hpp b/src/include/mallocMC/reservePoolPolicies/SimpleMalloc.hpp
@@ -1,6 +1,11 @@
 /*
   mallocMC: Memory Allocator for Many Core Architectures.
 
+  Copyright 2020 Helmholtz-Zentrum Dresden - Rossendorf,
+                 CERN
+
+  Author(s):  Bernhard Manfred Gruber
+
   Permission is hereby granted, free of charge, to any person obtaining a copy
   of this software and associated documentation files (the "Software"), to deal
   in the Software without restriction, including without limitation the rights

diff --git a/tests/verify_heap.cpp b/tests/verify_heap.cpp
@@ -99,23 +99,6 @@ auto main(int argc, char ** argv) -> int
 
     parse_cmdline(argc, argv, &heapInMB, &threads, &blocks, &machine_readable);
 
-    // TODO alpaka
-#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
-    int computeCapabilityMajor = 0;
-    cudaDeviceGetAttribute(
-        &computeCapabilityMajor, cudaDevAttrComputeCapabilityMajor, 0);
-    int computeCapabilityMinor = 0;
-    cudaDeviceGetAttribute(
-        &computeCapabilityMinor, cudaDevAttrComputeCapabilityMinor, 0);
-    if(computeCapabilityMajor < 2)
-    {
-        std::cerr << "Error: Compute Capability >= 2.0 required. (is ";
-        std::cerr << computeCapabilityMajor << "." << computeCapabilityMinor
-                  << ")\n";
-        return 1;
-    }
-#endif
-
     const auto correct
         = run_heap_verification(heapInMB, threads, blocks, machine_readable);
     if(!machine_readable || verbose)
@@ -663,9 +646,6 @@ auto run_heap_verification(
     const auto dev = alpaka::pltf::getDevByIdx<alpaka::pltf::Pltf<Device>>(0);
     auto queue = Queue{dev};
 
-    // cudaSetDeviceFlags(cudaDeviceMapHost); // TODO alpaka, and do we actually
-    // need this?
-
     const size_t heapSize = size_t(1024U * 1024U) * heapMB;
     const size_t slotSize = sizeof(allocElem_t) * ELEMS_PER_SLOT;
     const size_t nPointers = (heapSize + slotSize - 1) / slotSize;