Fix MSVC CUDA build: is_unsigned_v not available in cutlass::platform

include/cutlass/platform/platform.h

@@ -524,14 +524,17 @@ using std::is_trivially_copyable;
 
 #endif
 
-#if (CUTLASS_CXX17_OR_LATER)
+// Ensure is_unsigned is available (standard in <type_traits>).
+using CUTLASS_STL_NAMESPACE::is_unsigned;
 
-/// std::is_unsigned_v
-using CUTLASS_STL_NAMESPACE::is_integral_v;
-/// std::is_unsigned_v
-using CUTLASS_STL_NAMESPACE::is_unsigned_v;
+// Provide stable, C++11-compatible definitions for is_integral_v and is_unsigned_v.
+// This avoids issues on MSVC/CUDA where CUTLASS_CXX17_OR_LATER is true but the
+// underlying STL implementation does not expose the _v variants.
+template <typename T>
+constexpr bool is_integral_v = is_integral<T>::value;
 
-#endif
+template <typename T>
+constexpr bool is_unsigned_v = is_unsigned<T>::value;
 
 //-----------------------------------------------------------------------------
 // <utility>

test/unit/core/CMakeLists.txt

@@ -45,4 +45,5 @@ cutlass_test_unit_add_executable(
   numeric_conversion_subbyte.cu
   fast_numeric_conversion.cu
   functional.cu
+  platform_traits.cpp
   )

test/unit/core/platform_traits.cpp

@@ -0,0 +1,53 @@
+/***************************************************************************************************
+ * Copyright (c) 2017 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ * list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+ * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+ * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ **************************************************************************************************/
+
+#include "cutlass/platform/platform.h"
+
+#include "../common/cutlass_unit_test.h"
+
+TEST(PlatformTraits, IsIntegralV)
+{
+  // Compile-time checks to ensure is_integral_v is correctly defined.
+  static_assert(cutlass::platform::is_integral_v<int>,      "int is integral");
+  static_assert(cutlass::platform::is_integral_v<long>,     "long is integral");
+  static_assert(cutlass::platform::is_integral_v<unsigned int>, "unsigned int is integral");
+  static_assert(!cutlass::platform::is_integral_v<float>,   "float is not integral");
+  static_assert(!cutlass::platform::is_integral_v<double>,  "double is not integral");
+}
+
+TEST(PlatformTraits, IsUnsignedV)
+{
+  // Compile-time checks to ensure is_unsigned_v is correctly defined.
+  static_assert(cutlass::platform::is_unsigned_v<unsigned int>, "unsigned int is unsigned");
+  static_assert(cutlass::platform::is_unsigned_v<unsigned long>, "unsigned long is unsigned");
+  static_assert(!cutlass::platform::is_unsigned_v<int>, "int is not unsigned");
+  static_assert(!cutlass::platform::is_unsigned_v<float>, "float is not unsigned");
+}