[SYCLomatic] Support migration of cub::{SmVersion, SmVersionUncached} (…

…#2091) Signed-off-by: Wang, Yihan <[email protected]>
oneapi-src · Jun 26, 2024 · e6ee3c6 · e6ee3c6
1 parent dcd88af
commit e6ee3c6
Show file tree

Hide file tree

Showing 4 changed files with 48 additions and 4 deletions.
diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc
@@ -225,13 +225,13 @@ ENTRY_MEMBER_FUNCTION(cub::ChainedPolicy, cub::ChainedPolicy, Invoke, Invoke, fa
 ENTRY(cub::Debug, cub::Debug, false, NO_FLAG, P4, "Comment")
 ENTRY(cub::MaxSmOccupancy, cub::MaxSmOccupancy, false, NO_FLAG, P4, "Comment")
 ENTRY(cub::PtxVersion, cub::PtxVersion, true, NO_FLAG, P4, "Successful")
-ENTRY(cub::SmVersion, cub::SmVersion, false, NO_FLAG, P4, "Comment")
+ENTRY(cub::SmVersion, cub::SmVersion, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::CurrentDevice, cub::CurrentDevice, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::DeviceCount, cub::DeviceCount, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::DeviceCountUncached, cub::DeviceCountUncached, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::DeviceCountCachedValue, cub::DeviceCountCachedValue, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::PtxVersionUncached, cub::PtxVersionUncached, true, NO_FLAG, P4, "Successful")
-ENTRY(cub::SmVersionUncached, cub::SmVersionUncached, false, NO_FLAG, P4, "Comment")
+ENTRY(cub::SmVersionUncached, cub::SmVersionUncached, true, NO_FLAG, P4, "Successful")
 ENTRY(cub::SyncStream, cub::SyncStream, true, NO_FLAG, P4, "Successful")
 
 // Fancy iterators

diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp
@@ -189,7 +189,8 @@ void CubIntrinsicRule::registerMatcher(ast_matchers::MatchFinder &MF) {
               hasAnyName("IADD3", "SHR_ADD", "SHL_ADD", "BFE", "BFI", "LaneId",
                          "WarpId", "SyncStream", "CurrentDevice", "DeviceCount",
                          "DeviceCountUncached", "DeviceCountCachedValue",
-                         "PtxVersion", "PtxVersionUncached"),
+                         "PtxVersion", "PtxVersionUncached", "SmVersion",
+                         "SmVersionUncached"),
               hasAncestor(namespaceDecl(hasName("cub")))))))
           .bind("IntrinsicCall"),
       this);

diff --git a/clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp b/clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp
@@ -86,5 +86,29 @@ RewriterMap dpct::createUtilityFunctionsRewriterMap() {
                            LITERAL("DPCT_COMPATIBILITY_TEMP"))
       // cub::PtxVersionUncached
       ASSIGN_FACTORY_ENTRY("cub::PtxVersionUncached", ARG(0),
-                           LITERAL("DPCT_COMPATIBILITY_TEMP"))};
+                           LITERAL("DPCT_COMPATIBILITY_TEMP"))
+      // cub::SmVersion
+      ASSIGN_FACTORY_ENTRY(
+          "cub::SmVersion", ARG(0),
+          BO(BO_Add,
+             BO(BO_Mul,
+                CALL(MapNames::getDpctNamespace() + "get_major_version",
+                     makeDeviceStr()),
+                LITERAL("100")),
+             BO(BO_Mul,
+                CALL(MapNames::getDpctNamespace() + "get_minor_version",
+                     makeDeviceStr()),
+                LITERAL("10"))))
+      // cub::SmVersionUncached
+      ASSIGN_FACTORY_ENTRY(
+          "cub::SmVersionUncached", ARG(0),
+          BO(BO_Add,
+             BO(BO_Mul,
+                CALL(MapNames::getDpctNamespace() + "get_major_version",
+                     makeDeviceStr()),
+                LITERAL("100")),
+             BO(BO_Mul,
+                CALL(MapNames::getDpctNamespace() + "get_minor_version",
+                     makeDeviceStr()),
+                LITERAL("10"))))};
 }
diff --git a/clang/test/dpct/cub/intrinsic/sm_version.cu b/clang/test/dpct/cub/intrinsic/sm_version.cu
@@ -0,0 +1,19 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
+// RUN: dpct --format-range=none -in-root %S -out-root %T/intrinsic/sm_version %S/sm_version.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/intrinsic/sm_version/sm_version.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/intrinsic/sm_version/sm_version.dp.cpp -o %T/intrinsic/sm_version/sm_version.dp.o %}
+
+#include <cub/cub.cuh>
+
+void test() {
+  int a = 0;
+  // CHECK: a = dpct::get_major_version(dev_ct1) * 100 + dpct::get_minor_version(dev_ct1) * 10;
+  // CHECK-NEXT: a = dpct::get_major_version(dev_ct1) * 100 + dpct::get_minor_version(dev_ct1) * 10;
+  // CHECK-NEXT: a = dpct::get_major_version(dev_ct1) * 100 + dpct::get_minor_version(dev_ct1) * 10;
+  // CHECK-NEXT: a = dpct::get_major_version(dev_ct1) * 100 + dpct::get_minor_version(dev_ct1) * 10;
+  cub::SmVersion(a);
+  cub::SmVersion(a, 0);
+  cub::SmVersionUncached(a);
+  cub::SmVersionUncached(a, 0);
+}