Merge pull request #185 from NVIDIA/bugfix/test_platform_arch

Fix nv/target's platform test when compiling for newer arch
NVIDIA · Jul 15, 2021 · 11c5e75 · 11c5e75
2 parents 80fc78c + 555ec02
commit 11c5e75
Showing 1 changed file with 14 additions and 12 deletions.
diff --git a/.upstream-tests/test/cuda/test_platform.pass.cpp b/.upstream-tests/test/cuda/test_platform.pass.cpp
@@ -31,23 +31,25 @@ __host__ __device__ void test() {
   // This test ensures that the fallthrough cases are not invoked.
   // SM_80 would imply that SM_72 is available, yet it should not be expanded by the macro
   NV_DISPATCH_TARGET(
-    NV_PROVIDES_SM_80, (static_assert(arch_val == 800, "cuda arch expected 800");),
-    NV_PROVIDES_SM_75, (static_assert(arch_val == 750, "cuda arch expected 750");),
-    NV_PROVIDES_SM_72, (static_assert(arch_val == 720, "cuda arch expected 720");),
-    NV_PROVIDES_SM_70, (static_assert(arch_val == 700, "cuda arch expected 700");),
-    NV_PROVIDES_SM_62, (static_assert(arch_val == 620, "cuda arch expected 620");),
-    NV_PROVIDES_SM_61, (static_assert(arch_val == 610, "cuda arch expected 610");),
-    NV_PROVIDES_SM_60, (static_assert(arch_val == 600, "cuda arch expected 600");),
-    NV_PROVIDES_SM_53, (static_assert(arch_val == 530, "cuda arch expected 530");),
-    NV_PROVIDES_SM_52, (static_assert(arch_val == 520, "cuda arch expected 520");),
-    NV_PROVIDES_SM_50, (static_assert(arch_val == 500, "cuda arch expected 500");),
-    NV_PROVIDES_SM_37, (static_assert(arch_val == 370, "cuda arch expected 370");),
-    NV_PROVIDES_SM_35, (static_assert(arch_val == 350, "cuda arch expected 350");),
+    NV_PROVIDES_SM_86, (static_assert(arch_val >= 860, "cuda arch expected 860");),
+    NV_PROVIDES_SM_80, (static_assert(arch_val >= 800, "cuda arch expected 800");),
+    NV_PROVIDES_SM_75, (static_assert(arch_val >= 750, "cuda arch expected 750");),
+    NV_PROVIDES_SM_72, (static_assert(arch_val >= 720, "cuda arch expected 720");),
+    NV_PROVIDES_SM_70, (static_assert(arch_val >= 700, "cuda arch expected 700");),
+    NV_PROVIDES_SM_62, (static_assert(arch_val >= 620, "cuda arch expected 620");),
+    NV_PROVIDES_SM_61, (static_assert(arch_val >= 610, "cuda arch expected 610");),
+    NV_PROVIDES_SM_60, (static_assert(arch_val >= 600, "cuda arch expected 600");),
+    NV_PROVIDES_SM_53, (static_assert(arch_val >= 530, "cuda arch expected 530");),
+    NV_PROVIDES_SM_52, (static_assert(arch_val >= 520, "cuda arch expected 520");),
+    NV_PROVIDES_SM_50, (static_assert(arch_val >= 500, "cuda arch expected 500");),
+    NV_PROVIDES_SM_37, (static_assert(arch_val >= 370, "cuda arch expected 370");),
+    NV_PROVIDES_SM_35, (static_assert(arch_val >= 350, "cuda arch expected 350");),
     NV_IS_HOST,        (static_assert(arch_val == 0,   "cuda arch expected 0");)
   )
 
   // This test is simpler and ensures that only the value matched is invoked, but is roughly the same as the above
   NV_DISPATCH_TARGET(
+    NV_IS_EXACTLY_SM_86, (static_assert(arch_val == 860, "cuda arch expected 800");),
     NV_IS_EXACTLY_SM_80, (static_assert(arch_val == 800, "cuda arch expected 800");),
     NV_IS_EXACTLY_SM_75, (static_assert(arch_val == 750, "cuda arch expected 750");),
     NV_IS_EXACTLY_SM_72, (static_assert(arch_val == 720, "cuda arch expected 720");),