[SYCLomatic] Support query CUDA syntax sugar of --query-api-mapping. (#…

…2376) Signed-off-by: Tang, Jiajun [email protected]
oneapi-src · Sep 30, 2024 · c541cf0 · c541cf0
1 parent 234abeb
commit c541cf0
Show file tree

Hide file tree

Showing 10 changed files with 147 additions and 24 deletions.
diff --git a/clang/examples/DPCT/sugar/__constant__.cu b/clang/examples/DPCT/sugar/__constant__.cu
@@ -0,0 +1,3 @@
+// Start
+__constant__ int v;
+// End
diff --git a/clang/examples/DPCT/sugar/__device__.cu b/clang/examples/DPCT/sugar/__device__.cu
@@ -0,0 +1,4 @@
+// Start
+__device__ int v;
+__device__ void f() {}
+// End
diff --git a/clang/examples/DPCT/sugar/__global__.cu b/clang/examples/DPCT/sugar/__global__.cu
@@ -0,0 +1,3 @@
+// Start
+__global__ void f() {}
+// End
diff --git a/clang/examples/DPCT/sugar/__host__.cu b/clang/examples/DPCT/sugar/__host__.cu
@@ -0,0 +1,3 @@
+// Start
+__host__ void f() {}
+// End
diff --git a/clang/examples/DPCT/sugar/__managed__.cu b/clang/examples/DPCT/sugar/__managed__.cu
@@ -0,0 +1,3 @@
+// Start
+__managed__ int v;
+// End
diff --git a/clang/examples/DPCT/sugar/__shared__.cu b/clang/examples/DPCT/sugar/__shared__.cu
@@ -0,0 +1,3 @@
+// Start
+__global__ void f() { __shared__ int v; }
+// End
diff --git a/clang/examples/DPCT/sugar/kernel.cu b/clang/examples/DPCT/sugar/kernel.cu
@@ -0,0 +1,8 @@
+__global__ void f() {}
+
+void test() {
+  dim3 gridDim, blockDim;
+  // Start
+  f<<<gridDim, blockDim>>>();
+  // End
+}
diff --git a/clang/lib/DPCT/APIMapping/QueryAPIMapping.cpp b/clang/lib/DPCT/APIMapping/QueryAPIMapping.cpp
@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "QueryAPIMapping.h"
+#include "llvm/ADT/Twine.h"
 #include "llvm/Support/raw_ostream.h"
 #include <unordered_map>
 
@@ -27,34 +28,32 @@ void APIMapping::registerEntry(std::string Name, llvm::StringRef SourceCode) {
   const auto TargetIndex = EntryArray.size();
   EntryMap[Name] = TargetIndex; // Set the entry whether it exist or not.
   // Try to fuzz the original API name (only when the entry not exist):
-  // 1. Remove partial or all leading '_'.
-  // 2. For each name got by step 1, put 4 kind of fuzzed name into the map
+  // 1. Change "Name" to lower case. (Querying will change "Key" to lower too)
+  // 2. Remove partial or all suffix '_'.
+  std::transform(Name.begin(), Name.end(), Name.begin(), ::tolower);
+  while (Name.back() == '_') {
+    Name.erase(Name.end() - 1);
+    EntryMap.try_emplace(Name, TargetIndex);
+  }
+  const auto EmplaceWithAndWithoutSuffix = [TargetIndex](
+                                               llvm::StringRef Name,
+                                               llvm::StringRef Suffix) {
+    EntryMap.try_emplace(Name.str(), TargetIndex);
+    if (Name.take_back(Suffix.size()) == Suffix) {
+      EntryMap.try_emplace(Name.drop_back(Suffix.size()).str(), TargetIndex);
+    } else {
+      EntryMap.try_emplace(llvm::Twine(Name).concat(Suffix).str(), TargetIndex);
+    }
+  };
+  // 3. Remove partial or all leading '_'.
+  // 4. For each name got by step 1, put 2 kind of fuzzed name into the map
   // keys:
   //   (1) original name
   //   (2) remove or add Suffix "_v2"
-  //   (3) first char upper case name
-  //   (4) all char upper case name
-  //   (5) all char lower case name
-  for (int i = Name.find_first_not_of("_"); i >= 0; --i) {
-    auto TempName = Name;
-    std::string Suffix = "_v2";
-    if (TempName.size() > Suffix.length() &&
-        TempName.substr(TempName.size() - Suffix.length()) == Suffix) {
-      EntryMap.try_emplace(TempName.substr(0, TempName.size() - 3),
-                           TargetIndex);
-    } else {
-      EntryMap.try_emplace(TempName + Suffix, TargetIndex);
-    }
-    TempName[i] = std::toupper(TempName[i]);
-    EntryMap.try_emplace(TempName, TargetIndex);
-    std::transform(TempName.begin(), TempName.end(), TempName.begin(),
-                   ::toupper);
-    EntryMap.try_emplace(TempName, TargetIndex);
-    std::transform(TempName.begin(), TempName.end(), TempName.begin(),
-                   ::tolower);
-    EntryMap.try_emplace(TempName, TargetIndex);
+  EmplaceWithAndWithoutSuffix(Name, "_v2");
+  while (Name.front() == '_') {
     Name.erase(0, 1);
-    EntryMap.try_emplace(Name, TargetIndex);
+    EmplaceWithAndWithoutSuffix(Name, "_v2");
   }
   EntryArray.emplace_back(SourceCode);
 }
@@ -68,6 +67,10 @@ llvm::StringRef APIMapping::getAPISourceCode(std::string Key) {
   Key.erase(Key.find_last_not_of(" ") + 1);
   auto Iter = EntryMap.find(Key);
   if (Iter == EntryMap.end()) {
+    if (Key.find('<') != std::string::npos ||
+        Key.find('>') != std::string::npos) {
+      Key = "kernel";
+    }
     std::transform(Key.begin(), Key.end(), Key.begin(), ::tolower);
     Iter = EntryMap.find(Key);
   }

diff --git a/clang/test/dpct/query_api_mapping/sugar/test.cu b/clang/test/dpct/query_api_mapping/sugar/test.cu
@@ -0,0 +1,86 @@
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=kernel | FileCheck %s -check-prefix=KERNEL
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping="<<<>>>" | FileCheck %s -check-prefix=KERNEL
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping="<<<" | FileCheck %s -check-prefix=KERNEL
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=">>>" | FileCheck %s -check-prefix=KERNEL
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping="kernel<<<...>>>" | FileCheck %s -check-prefix=KERNEL
+
+// KERNEL: CUDA API:
+// KERNEL-NEXT:   f<<<gridDim, blockDim>>>();
+// KERNEL-NEXT: Is migrated to:
+// KERNEL-NEXT:   dpct::get_in_order_queue().parallel_for(
+// KERNEL-NEXT:     sycl::nd_range<3>(gridDim * blockDim, blockDim),
+// KERNEL-NEXT:     [=](sycl::nd_item<3> item_ct1) {
+// KERNEL-NEXT:       f();
+// KERNEL-NEXT:     });
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__constant__ | FileCheck %s -check-prefix=__CONSTANT__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__constant | FileCheck %s -check-prefix=__CONSTANT__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=constant | FileCheck %s -check-prefix=__CONSTANT__
+
+// __CONSTANT__: CUDA API:
+// __CONSTANT__-NEXT: __constant__ int v;
+// __CONSTANT__-NEXT: Is migrated to:
+// __CONSTANT__-NEXT: static dpct::constant_memory<int, 0> v;
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__device__ | FileCheck %s -check-prefix=__DEVICE__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__device | FileCheck %s -check-prefix=__DEVICE__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=device | FileCheck %s -check-prefix=__DEVICE__
+
+// __DEVICE__: CUDA API:
+// __DEVICE__-NEXT: __device__ int v;
+// __DEVICE__-NEXT: __device__ void f() {}
+// __DEVICE__-NEXT: Is migrated to:
+// __DEVICE__-NEXT: dpct::global_memory<int, 0> v;
+// __DEVICE__-NEXT: void f() {}
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__global__ | FileCheck %s -check-prefix=__GLOBAL__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__global | FileCheck %s -check-prefix=__GLOBAL__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=global | FileCheck %s -check-prefix=__GLOBAL__
+
+// __GLOBAL__: CUDA API:
+// __GLOBAL__-NEXT: __global__ void f() {}
+// __GLOBAL__-NEXT: Is migrated to:
+// __GLOBAL__-NEXT: void f() {}
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__host__ | FileCheck %s -check-prefix=__HOST__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__host | FileCheck %s -check-prefix=__HOST__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=host | FileCheck %s -check-prefix=__HOST__
+
+// __HOST__: CUDA API:
+// __HOST__-NEXT: __host__ void f() {}
+// __HOST__-NEXT: Is migrated to:
+// __HOST__-NEXT: void f() {}
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__managed__ | FileCheck %s -check-prefix=__MANAGED__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__managed | FileCheck %s -check-prefix=__MANAGED__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=managed | FileCheck %s -check-prefix=__MANAGED__
+
+// __MANAGED__: CUDA API:
+// __MANAGED__-NEXT: __managed__ int v;
+// __MANAGED__-NEXT: Is migrated to:
+// __MANAGED__-NEXT: dpct::shared_memory<int, 0> v;
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__shared__ | FileCheck %s -check-prefix=__SHARED__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=__shared | FileCheck %s -check-prefix=__SHARED__
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=shared | FileCheck %s -check-prefix=__SHARED__
+
+// __SHARED__: CUDA API:
+// __SHARED__-NEXT: __global__ void f() { __shared__ int v; }
+// __SHARED__-NEXT: Is migrated to:
+// __SHARED__-NEXT: void f(int &v) { }
diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu
@@ -45,6 +45,7 @@
 // CHECK-NEXT: __byte_perm
 // CHECK-NEXT: __clz
 // CHECK-NEXT: __clzll
+// CHECK-NEXT: __constant__
 // CHECK-NEXT: __cosf
 // CHECK-NEXT: __dadd_rd
 // CHECK-NEXT: __dadd_rn
@@ -54,6 +55,7 @@
 // CHECK-NEXT: __ddiv_rn
 // CHECK-NEXT: __ddiv_ru
 // CHECK-NEXT: __ddiv_rz
+// CHECK-NEXT: __device__
 // CHECK-NEXT: __dmul_rd
 // CHECK-NEXT: __dmul_rn
 // CHECK-NEXT: __dmul_ru
@@ -167,6 +169,7 @@
 // CHECK-NEXT: __fsub_rn
 // CHECK-NEXT: __fsub_ru
 // CHECK-NEXT: __fsub_rz
+// CHECK-NEXT: __global__
 // CHECK-NEXT: __h2div
 // CHECK-NEXT: __habs
 // CHECK-NEXT: __habs2
@@ -290,6 +293,7 @@
 // CHECK-NEXT: __hneu
 // CHECK-NEXT: __hneu2
 // CHECK-NEXT: __hneu2_mask
+// CHECK-NEXT: __host__
 // CHECK-NEXT: __hsub
 // CHECK-NEXT: __hsub2
 // CHECK-NEXT: __hsub2_rn
@@ -344,6 +348,7 @@
 // CHECK-NEXT: __lowhigh2highlow
 // CHECK-NEXT: __lows2bfloat162
 // CHECK-NEXT: __lows2half2
+// CHECK-NEXT: __managed__
 // CHECK-NEXT: __match_all_sync
 // CHECK-NEXT: __match_any_sync
 // CHECK-NEXT: __mul24
@@ -362,6 +367,7 @@
 // CHECK-NEXT: __rhadd
 // CHECK-NEXT: __sad
 // CHECK-NEXT: __saturatef
+// CHECK-NEXT: __shared__
 // CHECK-NEXT: __shfl
 // CHECK-NEXT: __shfl_down
 // CHECK-NEXT: __shfl_down_sync
@@ -1832,6 +1838,7 @@
 // CHECK-NEXT: j1f
 // CHECK-NEXT: jn
 // CHECK-NEXT: jnf
+// CHECK-NEXT: kernel
 // CHECK-NEXT: ldexp
 // CHECK-NEXT: ldexpf
 // CHECK-NEXT: lgamma