 /**
  * \brief Retrieves the PTX version that will be used on \p device (major * 100 + minor * 10).
  */
-__host__ __forceinline__ cudaError_t PtxVersionUncached(int &ptx_version, int device)
+__host__ __forceinline__ cudaError_t PtxVersionUncached(int& ptx_version, int device)
 {
     SwitchDevice sd(device);
-    // Neither the single argument version of `PtxVersionUncached` or this
-    // overload passes the `cudaError_t` through `CubDebug` because failure
-    // is not a hard error here. We may be querying a device that we do not
-    // have code for but never use.
     return PtxVersionUncached(ptx_version);
 }
 
+#if __cplusplus >= 201103L // C++11 and later.
+template <typename Tag>
+__host__ __forceinline__ PerDeviceAttributeCache& GetPerDeviceAttributeCache()
+{
+    // C++11 guarantees that initialization of static locals is thread safe.
+    static PerDeviceAttributeCache cache;
+    return cache;
+}
+
+struct PtxVersionCacheTag {};
+struct SmVersionCacheTag {};
+#endif
+
 /**
  * \brief Retrieves the PTX version that will be used on \p device (major * 100 + minor * 10).
  *
  * \note This function may cache the result internally.
- */
-__host__ __forceinline__ cudaError_t PtxVersion(int &ptx_version, int device)
+ *
+ * \note This function is thread safe.
+ */
+__host__ __forceinline__ cudaError_t PtxVersion(int& ptx_version, int device)
 {
 #if __cplusplus >= 201103L // C++11 and later.
 
-    using FunctionPointer = cudaError_t(*)(int &, int);
-    FunctionPointer fun_ptr = PtxVersionUncached;
-
-    // C++11 guarantees that initialization of static locals is thread safe.
-    static const PerDeviceAttributeCache cache(fun_ptr);
-
-    if (!CubDebug(cache.error[device]))
-        ptx_version = cache.attribute[device];
-
-    return cache.error[device];
+    auto const payload = GetPerDeviceAttributeCache<PtxVersionCacheTag>()(
+      // If this call fails, then we get the error code back in the payload,
+      // which we check with `CubDebug` below.
+      [=] (int& pv) { return PtxVersionUncached(pv, device); },
+      device);
+
+    if (!CubDebug(payload.error))
+        ptx_version = payload.attribute;
+
+    return payload.error;
 
 #else // Pre C++11.
 
-    return CubDebug(PtxVersionUncached(ptx_version, device));
+    return PtxVersionUncached(ptx_version, device);
 
 #endif
 }

__host__ __device__ __forceinline__
cudaError_t AliasTemporaries(
    void    *d_temp_storage,                    ///< [in] %Device-accessible allocation of temporary storage.  When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done.
    size_t  &temp_storage_bytes,                ///< [in,out] Size in bytes of \t d_temp_storage allocation
    void*   (&allocations)[ALLOCATIONS],        ///< [in,out] Pointers to device allocations needed
    size_t  (&allocation_sizes)[ALLOCATIONS])   ///< [in] Sizes in bytes of device allocations needed
{

   /// Specializes and dispatches op in accordance to the first policy in the chain of adequate PTX version
   template <typename FunctorT>
   CUB_RUNTIME_FUNCTION __forceinline__
   static cudaError_t Invoke(int ptx_version, FunctorT& op)
   {
       if (ptx_version < PTX_VERSION) {
           return PrevPolicyT::Invoke(ptx_version, op);
