Implement cache of device properties

include/alpaka/dev/DevCpu.hpp

@@ -89,6 +89,12 @@ namespace alpaka
             return 0;
         }
 
+        friend struct trait::GetName<DevCpu>;
+        friend struct trait::GetMemBytes<DevCpu>;
+        friend struct trait::GetFreeMemBytes<DevCpu>;
+        friend struct trait::GetWarpSizes<DevCpu>;
+        friend struct trait::GetPreferredWarpSize<DevCpu>;
+
     private:
         std::shared_ptr<cpu::detail::DevCpuImpl> m_spDevCpuImpl;
     };
@@ -99,19 +105,35 @@ namespace alpaka
         template<>
         struct GetName<DevCpu>
         {
-            ALPAKA_FN_HOST static auto getName(DevCpu const& /* dev */) -> std::string
+            ALPAKA_FN_HOST static auto getName(DevCpu const& dev) -> std::string
             {
-                return cpu::detail::getCpuName();
+                auto& name = dev.m_spDevCpuImpl->deviceProperties().name;
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_spDevCpuImpl->mutex());
+                    if(!name.has_value())
+                    {
+                        name = cpu::detail::getCpuName();
+                    }
+                }
+                return name.value();
             }
         };
 
         //! The CPU device available memory get trait specialization.
         template<>
         struct GetMemBytes<DevCpu>
         {
-            ALPAKA_FN_HOST static auto getMemBytes(DevCpu const& /* dev */) -> std::size_t
+            ALPAKA_FN_HOST static auto getMemBytes(DevCpu const& dev) -> std::size_t
             {
-                return cpu::detail::getTotalGlobalMemSizeBytes();
+                auto& totalGlobalMem = dev.m_spDevCpuImpl->deviceProperties().totalGlobalMem;
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_spDevCpuImpl->mutex());
+                    if(!totalGlobalMem.has_value())
+                    {
+                        totalGlobalMem = cpu::detail::getTotalGlobalMemSizeBytes();
+                    }
+                }
+                return totalGlobalMem.value();
             }
         };
 

include/alpaka/dev/DevGenericSycl.hpp

@@ -9,6 +9,7 @@
 #include "alpaka/core/Common.hpp"
 #include "alpaka/core/Sycl.hpp"
 #include "alpaka/dev/Traits.hpp"
+#include "alpaka/dev/common/DeviceProperties.hpp"
 #include "alpaka/mem/buf/Traits.hpp"
 #include "alpaka/platform/Traits.hpp"
 #include "alpaka/queue/Properties.hpp"
@@ -112,10 +113,21 @@ namespace alpaka
                 return m_context;
             }
 
+            std::shared_mutex& mutex()
+            {
+                return m_mutex;
+            }
+
+            alpaka::DeviceProperties& deviceProperties()
+            {
+                return m_deviceProperties;
+            }
+
         private:
             sycl::device m_device;
             sycl::context m_context;
             std::vector<std::weak_ptr<QueueGenericSyclImpl>> m_queues;
+            alpaka::DeviceProperties m_deviceProperties;
             std::shared_mutex mutable m_mutex;
         };
     } // namespace detail
@@ -160,8 +172,16 @@ namespace alpaka
         {
             static auto getName(DevGenericSycl<TTag> const& dev) -> std::string
             {
-                auto const device = dev.getNativeHandle().first;
-                return device.template get_info<sycl::info::device::name>();
+                auto& name = dev.m_impl->deviceProperties().name;
+                {
+                    std::lock_guard<std::shared_mutex> lock(dev.m_impl->mutex());
+                    if(!name.has_value())
+                    {
+                        auto const device = dev.getNativeHandle().first;
+                        name = device.template get_info<sycl::info::device::name>();
+                    }
+                }
+                return name.value();
             }
         };
 
@@ -171,8 +191,16 @@ namespace alpaka
         {
             static auto getMemBytes(DevGenericSycl<TTag> const& dev) -> std::size_t
             {
-                auto const device = dev.getNativeHandle().first;
-                return device.template get_info<sycl::info::device::global_mem_size>();
+                auto& totalGlobalMem = dev.m_impl->deviceProperties().totalGlobalMem;
+                {
+                    std::lock_guard<std::shared_mutex> lock(dev.m_impl->mutex());
+                    if(!totalGlobalMem.has_value())
+                    {
+                        auto const device = dev.getNativeHandle().first;
+                        totalGlobalMem = device.template get_info<sycl::info::device::global_mem_size>();
+                    }
+                }
+                return totalGlobalMem.value();
             }
         };
 
@@ -195,15 +223,25 @@ namespace alpaka
         {
             static auto getWarpSizes(DevGenericSycl<TTag> const& dev) -> std::vector<std::size_t>
             {
-                auto const device = dev.getNativeHandle().first;
-                std::vector<std::size_t> warp_sizes = device.template get_info<sycl::info::device::sub_group_sizes>();
-                // The CPU runtime supports a sub-group size of 64, but the SYCL implementation currently does not
-                auto find64 = std::find(warp_sizes.begin(), warp_sizes.end(), 64);
-                if(find64 != warp_sizes.end())
-                    warp_sizes.erase(find64);
-                // Sort the warp sizes in decreasing order
-                std::sort(warp_sizes.begin(), warp_sizes.end(), std::greater<>{});
-                return warp_sizes;
+                auto& warpSizes = dev.m_impl->deviceProperties().warpSizes;
+                {
+                    std::lock_guard<std::shared_mutex> lock(dev.m_impl->mutex());
+                    if(!warpSizes.has_value())
+                    {
+                        auto const device = dev.getNativeHandle().first;
+                        std::vector<std::size_t> warp_sizes
+                            = device.template get_info<sycl::info::device::sub_group_sizes>();
+                        // The CPU runtime supports a sub-group size of 64, but the SYCL implementation currently does
+                        // not
+                        auto find64 = std::find(warp_sizes.begin(), warp_sizes.end(), 64);
+                        if(find64 != warp_sizes.end())
+                            warp_sizes.erase(find64);
+                        // Sort the warp sizes in decreasing order
+                        std::sort(warp_sizes.begin(), warp_sizes.end(), std::greater<>{});
+                        warpSizes = std::move(warp_sizes);
+                    }
+                }
+                return warpSizes.value();
             }
         };
 
@@ -213,6 +251,14 @@ namespace alpaka
         {
             static auto getPreferredWarpSize(DevGenericSycl<TTag> const& dev) -> std::size_t
             {
+                auto& warpSizes = dev.m_impl->deviceProperties().warpSizes;
+                {
+                    std::lock_guard<std::shared_mutex> lock(dev.m_impl->mutex());
+                    if(warpSizes.has_value())
+                    {
+                        return warpSizes.value().front();
+                    }
+                }
                 return GetWarpSizes<DevGenericSycl<TTag>>::getWarpSizes(dev).front();
             }
         };

include/alpaka/dev/DevUniformCudaHipRt.hpp

@@ -10,6 +10,7 @@
 #include "alpaka/core/Hip.hpp"
 #include "alpaka/core/Interface.hpp"
 #include "alpaka/dev/Traits.hpp"
+#include "alpaka/dev/common/DeviceProperties.hpp"
 #include "alpaka/dev/common/QueueRegistry.hpp"
 #include "alpaka/mem/buf/Traits.hpp"
 #include "alpaka/platform/Traits.hpp"
@@ -20,6 +21,7 @@
 #include "alpaka/wait/Traits.hpp"
 
 #include <cstddef>
+#include <mutex>
 #include <string>
 #include <vector>
 
@@ -94,6 +96,12 @@ namespace alpaka
             m_QueueRegistry->registerQueue(spQueue);
         }
 
+        friend struct trait::GetName<DevUniformCudaHipRt<TApi>>;
+        friend struct trait::GetMemBytes<DevUniformCudaHipRt<TApi>>;
+        friend struct trait::GetFreeMemBytes<DevUniformCudaHipRt<TApi>>;
+        friend struct trait::GetWarpSizes<DevUniformCudaHipRt<TApi>>;
+        friend struct trait::GetPreferredWarpSize<DevUniformCudaHipRt<TApi>>;
+
     private:
         DevUniformCudaHipRt(int iDevice)
             : m_iDevice(iDevice)
@@ -114,12 +122,20 @@ namespace alpaka
         {
             ALPAKA_FN_HOST static auto getName(DevUniformCudaHipRt<TApi> const& dev) -> std::string
             {
-                // There is cuda/hip-DeviceGetAttribute as faster alternative to cuda/hip-GetDeviceProperties to get a
-                // single device property but it has no option to get the name
-                typename TApi::DeviceProp_t devProp;
-                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::getDeviceProperties(&devProp, dev.getNativeHandle()));
-
-                return std::string(devProp.name);
+                auto& name = dev.m_QueueRegistry->deviceProperties().name;
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_QueueRegistry->mutex());
+                    if(!name.has_value())
+                    {
+                        // There is cuda/hip-DeviceGetAttribute as faster alternative to cuda/hip-GetDeviceProperties
+                        // to get a single device property but it has no option to get the name
+                        typename TApi::DeviceProp_t devProp;
+                        ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::getDeviceProperties(&devProp, dev.getNativeHandle()));
+                        name = std::string(devProp.name);
+                    }
+                }
+
+                return name.value();
             }
         };
 
@@ -129,15 +145,24 @@ namespace alpaka
         {
             ALPAKA_FN_HOST static auto getMemBytes(DevUniformCudaHipRt<TApi> const& dev) -> std::size_t
             {
-                // Set the current device to wait for.
-                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::setDevice(dev.getNativeHandle()));
+                auto& totalGlobalMem = dev.m_QueueRegistry->deviceProperties().totalGlobalMem;
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_QueueRegistry->mutex());
+                    if(!totalGlobalMem.has_value())
+                    {
+                        // Set the current device to wait for.
+                        ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::setDevice(dev.getNativeHandle()));
 
-                std::size_t freeInternal(0u);
-                std::size_t totalInternal(0u);
+                        std::size_t freeInternal(0u);
+                        std::size_t totalInternal(0u);
+
+                        ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::memGetInfo(&freeInternal, &totalInternal));
 
-                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::memGetInfo(&freeInternal, &totalInternal));
+                        totalGlobalMem = totalInternal;
+                    }
+                }
 
-                return totalInternal;
+                return totalGlobalMem.value();
             }
         };
 
@@ -147,13 +172,22 @@ namespace alpaka
         {
             ALPAKA_FN_HOST static auto getFreeMemBytes(DevUniformCudaHipRt<TApi> const& dev) -> std::size_t
             {
-                // Set the current device to wait for.
-                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::setDevice(dev.getNativeHandle()));
-
+                auto& totalGlobalMem = dev.m_QueueRegistry->deviceProperties().totalGlobalMem;
                 std::size_t freeInternal(0u);
-                std::size_t totalInternal(0u);
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_QueueRegistry->mutex());
+                    if(!totalGlobalMem.has_value())
+                    {
+                        // Set the current device to wait for.
+                        ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::setDevice(dev.getNativeHandle()));
 
-                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::memGetInfo(&freeInternal, &totalInternal));
+                        std::size_t totalInternal(0u);
+
+                        ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::memGetInfo(&freeInternal, &totalInternal));
+
+                        totalGlobalMem = totalInternal;
+                    }
+                }
 
                 return freeInternal;
             }
@@ -165,7 +199,28 @@ namespace alpaka
         {
             ALPAKA_FN_HOST static auto getWarpSizes(DevUniformCudaHipRt<TApi> const& dev) -> std::vector<std::size_t>
             {
-                return {GetPreferredWarpSize<DevUniformCudaHipRt<TApi>>::getPreferredWarpSize(dev)};
+                auto& warpSizes = dev.m_QueueRegistry->deviceProperties().warpSizes;
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_QueueRegistry->mutex());
+                    if(!warpSizes.has_value())
+                    {
+                        if(dev.m_QueueRegistry->deviceProperties().preferredWarpSize.has_value())
+                        {
+                            warpSizes = std::vector<std::size_t>{
+                                dev.m_QueueRegistry->deviceProperties().preferredWarpSize.value()};
+                        }
+                        else
+                        {
+                            int warpSize = 0;
+                            ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::deviceGetAttribute(
+                                &warpSize,
+                                TApi::deviceAttributeWarpSize,
+                                dev.getNativeHandle()));
+                            warpSizes = std::vector<std::size_t>{static_cast<std::size_t>(warpSize)};
+                        }
+                    }
+                }
+                return warpSizes.value();
             }
         };
 
@@ -175,11 +230,20 @@ namespace alpaka
         {
             ALPAKA_FN_HOST static auto getPreferredWarpSize(DevUniformCudaHipRt<TApi> const& dev) -> std::size_t
             {
-                int warpSize = 0;
-
-                ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(
-                    TApi::deviceGetAttribute(&warpSize, TApi::deviceAttributeWarpSize, dev.getNativeHandle()));
-                return static_cast<std::size_t>(warpSize);
+                auto& preferredWarpSize = dev.m_QueueRegistry->deviceProperties().preferredWarpSize;
+                {
+                    std::lock_guard<std::mutex> lock(dev.m_QueueRegistry->mutex());
+                    if(!preferredWarpSize.has_value())
+                    {
+                        int warpSize = 0;
+
+                        ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(
+                            TApi::deviceGetAttribute(&warpSize, TApi::deviceAttributeWarpSize, dev.getNativeHandle()));
+                        preferredWarpSize = static_cast<std::size_t>(warpSize);
+                    }
+                }
+
+                return preferredWarpSize.value();
             }
         };
 

include/alpaka/dev/common/DeviceProperties.hpp

@@ -0,0 +1,37 @@
+
+#pragma once
+
+#include "alpaka/dev/Traits.hpp"
+
+#include <optional>
+#include <string>
+#include <vector>
+
+namespace alpaka
+{
+    class DeviceProperties
+    {
+        std::optional<std::string> name;
+        std::optional<std::size_t> totalGlobalMem;
+        std::optional<std::vector<std::size_t>> warpSizes;
+        std::optional<std::size_t> preferredWarpSize;
+
+    public:
+        DeviceProperties() = default;
+
+        template<typename TDev, typename TSfinae>
+        friend struct trait::GetName;
+
+        template<typename TDev, typename TSfinae>
+        friend struct trait::GetMemBytes;
+
+        template<typename TDev, typename TSfinae>
+        friend struct trait::GetFreeMemBytes;
+
+        template<typename TDev, typename TSfinae>
+        friend struct trait::GetWarpSizes;
+
+        template<typename TDev, typename TSfinae>
+        friend struct trait::GetPreferredWarpSize;
+    };
+} // namespace alpaka