Add m_cooperativeLaunch device prop and runtime check for CG support …

…for CUDA/HIP
alpaka-group · Jul 29, 2024 · 565c944 · 565c944
1 parent 9614d9c
commit 565c944
Show file tree

Hide file tree

Showing 12 changed files with 62 additions and 10 deletions.
diff --git a/include/alpaka/acc/AccCpuOmp2Blocks.hpp b/include/alpaka/acc/AccCpuOmp2Blocks.hpp
@@ -151,7 +151,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         static_cast<size_t>(AccCpuOmp2Blocks<TDim, TIdx>::staticAllocBytes()),
                         // m_globalMemSizeBytes
-                        getMemBytes(dev)};
+                        getMemBytes(dev),
+                        //m_cooperativeLaunch
+                        std::false_type()};
             }
         };
 

diff --git a/include/alpaka/acc/AccCpuOmp2Threads.hpp b/include/alpaka/acc/AccCpuOmp2Threads.hpp
@@ -162,7 +162,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         memBytes,
                         // m_globalMemSizeBytes
-                        memBytes};
+                        memBytes,
+                        //m_cooperativeLaunch
+                        std::false_type()};
             }
         };
 

diff --git a/include/alpaka/acc/AccCpuSerial.hpp b/include/alpaka/acc/AccCpuSerial.hpp
@@ -145,7 +145,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         static_cast<size_t>(AccCpuSerial<TDim, TIdx>::staticAllocBytes()),
                         // m_globalMemSizeBytes
-                        getMemBytes(dev)};
+                        getMemBytes(dev),
+                        //m_cooperativeLaunch
+                        std::false_type()};
             }
         };
 

diff --git a/include/alpaka/acc/AccCpuTbbBlocks.hpp b/include/alpaka/acc/AccCpuTbbBlocks.hpp
@@ -143,7 +143,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         static_cast<size_t>(AccCpuTbbBlocks<TDim, TIdx>::staticAllocBytes()),
                         // m_globalMemSizeBytes
-                        getMemBytes(dev)};
+                        getMemBytes(dev),
+                        //m_cooperativeLaunch
+                        std::false_type()};
             }
         };
 

diff --git a/include/alpaka/acc/AccCpuThreads.hpp b/include/alpaka/acc/AccCpuThreads.hpp
@@ -170,7 +170,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         memBytes,
                         // m_globalMemSizeBytes
-                        memBytes};
+                        memBytes,
+                        //m_cooperativeLaunch
+                        std::false_type()};
             }
         };
 

diff --git a/include/alpaka/acc/AccDevProps.hpp b/include/alpaka/acc/AccDevProps.hpp
@@ -30,5 +30,6 @@ namespace alpaka
         TIdx m_threadElemCountMax; //!< The maximum number of elements in a threads.
         size_t m_sharedMemSizeBytes; //!< The size of shared memory per block
         size_t m_globalMemSizeBytes; //!< The size of global memory
+        bool m_cooperativeLaunch; //!< The support for launch of cooperative kernels
     };
 } // namespace alpaka
diff --git a/include/alpaka/acc/AccGenericSycl.hpp b/include/alpaka/acc/AccGenericSycl.hpp
@@ -153,7 +153,9 @@ namespace alpaka::trait
                     // m_sharedMemSizeBytes
                     device.template get_info<sycl::info::device::local_mem_size>(),
                     // m_globalMemSizeBytes
-                    getMemBytes(dev)};
+                    getMemBytes(dev),
+                    //m_cooperativeLaunch
+                    std::false_type()};
         }
     };
 

diff --git a/include/alpaka/acc/AccGpuUniformCudaHipRt.hpp b/include/alpaka/acc/AccGpuUniformCudaHipRt.hpp
@@ -163,6 +163,12 @@ namespace alpaka
                     TApi::deviceAttributeMaxSharedMemoryPerBlock,
                     dev.getNativeHandle()));
 
+                int cooperativeLaunch = {};
+                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::deviceGetAttribute(
+                    &cooperativeLaunch,
+                    TApi::deviceAttributeCooperativeLaunch,
+                    dev.getNativeHandle()));
+
                 return {// m_multiProcessorCount
                         alpaka::core::clipCast<TIdx>(multiProcessorCount),
                         // m_gridBlockExtentMax
@@ -186,7 +192,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         static_cast<size_t>(sharedMemSizeBytes),
                         // m_globalMemSizeBytes
-                        getMemBytes(dev)};
+                        getMemBytes(dev),
+                        //m_cooperativeLaunch
+                        static_cast<bool>(cooperativeLaunch)};
 
 #    else
                 typename TApi::DeviceProp_t properties;
@@ -215,7 +223,9 @@ namespace alpaka
                         // m_sharedMemSizeBytes
                         static_cast<size_t>(properties.sharedMemPerBlock),
                         // m_globalMemSizeBytes
-                        getMemBytes(dev)};
+                        getMemBytes(dev),
+                        //m_cooperativeLaunch
+                        static_cast<bool>(properties.cooperativeLaunch)};
 #    endif
             }
         };

diff --git a/include/alpaka/core/ApiCudaRt.hpp b/include/alpaka/core/ApiCudaRt.hpp
@@ -76,6 +76,7 @@ namespace alpaka
         static constexpr DeviceAttr_t deviceAttributeMaxThreadsPerBlock = ::cudaDevAttrMaxThreadsPerBlock;
         static constexpr DeviceAttr_t deviceAttributeMultiprocessorCount = ::cudaDevAttrMultiProcessorCount;
         static constexpr DeviceAttr_t deviceAttributeWarpSize = ::cudaDevAttrWarpSize;
+        static constexpr DeviceAttr_t deviceAttributeCooperativeLaunch = ::cudaDevAttrCooperativeLaunch;
 
         static constexpr Limit_t limitPrintfFifoSize = ::cudaLimitPrintfFifoSize;
         static constexpr Limit_t limitMallocHeapSize = ::cudaLimitMallocHeapSize;

diff --git a/include/alpaka/core/ApiHipRt.hpp b/include/alpaka/core/ApiHipRt.hpp
@@ -79,6 +79,7 @@ namespace alpaka
         static constexpr DeviceAttr_t deviceAttributeMaxThreadsPerBlock = ::hipDeviceAttributeMaxThreadsPerBlock;
         static constexpr DeviceAttr_t deviceAttributeMultiprocessorCount = ::hipDeviceAttributeMultiprocessorCount;
         static constexpr DeviceAttr_t deviceAttributeWarpSize = ::hipDeviceAttributeWarpSize;
+        static constexpr DeviceAttr_t deviceAttributeCooperativeLaunch = ::hipDeviceAttributeCooperativeLaunch;
 
 #    if HIP_VERSION >= 40'500'000
         static constexpr Limit_t limitPrintfFifoSize = ::hipLimitPrintfFifoSize;

diff --git a/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp b/include/alpaka/kernel/TaskKernelGpuUniformCudaHipRt.hpp
@@ -204,6 +204,18 @@ namespace alpaka
                 // TApi::deviceGetLimit(&printfFifoSize, TApi::limitPrintfFifoSize);
                 // std::cout << __func__ << "INFO: printfFifoSize: " << printfFifoSize << std::endl;
 #        endif
+
+#        if ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL
+                // This checks if the device supports cooperative kernel launch
+                if constexpr (TCooperative)
+                {
+                    if(!trait::GetAccDevProps<TAcc>::getAccDevProps(getDev(queue)).m_cooperativeLaunch)
+                    {
+                        throw std::runtime_error("This accelerator doesn't support cooperative gropus functionality!");
+                    } else std::cout << "This accelerator supports cooperative gropus functionality!";
+                }
+#        endif
+
                 auto const gridBlockExtent = getWorkDiv<Grid, Blocks>(task);
                 auto const blockThreadExtent = getWorkDiv<Block, Threads>(task);
                 auto const threadElemExtent = getWorkDiv<Thread, Elems>(task);
@@ -361,6 +373,18 @@ namespace alpaka
                 // TApi::deviceGetLimit(&printfFifoSize, TApi::limitPrintfFifoSize);
                 // std::cout << __func__ << "INFO: printfFifoSize: " << printfFifoSize << std::endl;
 #        endif
+
+#        if ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL
+                // This checks if the device supports cooperative kernel launch
+                if constexpr (TCooperative)
+                {
+                    if(!trait::GetAccDevProps<TAcc>::getAccDevProps(getDev(queue)).m_cooperativeLaunch)
+                    {
+                        throw std::runtime_error("This accelerator doesn't support cooperative groups functionality!");
+                    }
+                }
+#        endif
+
                 auto const gridBlockExtent = getWorkDiv<Grid, Blocks>(task);
                 auto const blockThreadExtent = getWorkDiv<Block, Threads>(task);
                 auto const threadElemExtent = getWorkDiv<Thread, Elems>(task);

diff --git a/test/unit/acc/src/AccDevPropsTest.cpp b/test/unit/acc/src/AccDevPropsTest.cpp
@@ -29,7 +29,7 @@ TEMPLATE_LIST_TEST_CASE("getAccDevProps", "[acc]", alpaka::test::TestAccs)
 
 TEST_CASE("AccDevProps.aggregate_init", "[acc]")
 {
-    auto const props = alpaka::AccDevProps<alpaka::DimInt<1>, int>{1, {2}, 3, {4}, 5, {6}, 7, 8, 9};
+    auto const props = alpaka::AccDevProps<alpaka::DimInt<1>, int>{1, {2}, 3, {4}, 5, {6}, 7, 8, 9, std::true_type()};
 
     CHECK(props.m_multiProcessorCount == 1);
     CHECK(props.m_gridBlockExtentMax == alpaka::Vec{2});
@@ -40,6 +40,7 @@ TEST_CASE("AccDevProps.aggregate_init", "[acc]")
     CHECK(props.m_threadElemCountMax == 7);
     CHECK(props.m_sharedMemSizeBytes == 8);
     CHECK(props.m_globalMemSizeBytes == 9);
+    CHECK(props.m_cooperativeLaunch == true);
 }
 
 #ifdef __cpp_designated_initializers
@@ -54,7 +55,8 @@ TEST_CASE("AccDevProps.designated_initializers", "[acc]")
         .m_threadElemExtentMax = {60},
         .m_threadElemCountMax = 70,
         .m_sharedMemSizeBytes = 80,
-        .m_globalMemSizeBytes = 90};
+        .m_globalMemSizeBytes = 90
+        .m_cooperativeLaunch = std::false_type()};
 
     CHECK(props.m_multiProcessorCount == 10);
     CHECK(props.m_gridBlockExtentMax == alpaka::Vec{20});
@@ -65,5 +67,6 @@ TEST_CASE("AccDevProps.designated_initializers", "[acc]")
     CHECK(props.m_threadElemCountMax == 70);
     CHECK(props.m_sharedMemSizeBytes == 80);
     CHECK(props.m_globalMemSizeBytes == 90);
+    CHECK(props.m_cooperativeLaunch == false);
 }
 #endif