From c541cf0b58121cdecbad6942676c0c502947f1f8 Mon Sep 17 00:00:00 2001 From: Tang Jiajun Date: Mon, 30 Sep 2024 09:34:21 +0800 Subject: [PATCH] [SYCLomatic] Support query CUDA syntax sugar of --query-api-mapping. (#2376) Signed-off-by: Tang, Jiajun jiajun.tang@intel.com --- clang/examples/DPCT/sugar/__constant__.cu | 3 + clang/examples/DPCT/sugar/__device__.cu | 4 + clang/examples/DPCT/sugar/__global__.cu | 3 + clang/examples/DPCT/sugar/__host__.cu | 3 + clang/examples/DPCT/sugar/__managed__.cu | 3 + clang/examples/DPCT/sugar/__shared__.cu | 3 + clang/examples/DPCT/sugar/kernel.cu | 8 ++ clang/lib/DPCT/APIMapping/QueryAPIMapping.cpp | 51 +++++------ .../test/dpct/query_api_mapping/sugar/test.cu | 86 +++++++++++++++++++ clang/test/dpct/query_api_mapping/test_all.cu | 7 ++ 10 files changed, 147 insertions(+), 24 deletions(-) create mode 100644 clang/examples/DPCT/sugar/__constant__.cu create mode 100644 clang/examples/DPCT/sugar/__device__.cu create mode 100644 clang/examples/DPCT/sugar/__global__.cu create mode 100644 clang/examples/DPCT/sugar/__host__.cu create mode 100644 clang/examples/DPCT/sugar/__managed__.cu create mode 100644 clang/examples/DPCT/sugar/__shared__.cu create mode 100644 clang/examples/DPCT/sugar/kernel.cu create mode 100644 clang/test/dpct/query_api_mapping/sugar/test.cu diff --git a/clang/examples/DPCT/sugar/__constant__.cu b/clang/examples/DPCT/sugar/__constant__.cu new file mode 100644 index 000000000000..579a1d383953 --- /dev/null +++ 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 new file mode 100644 index 000000000000..8b2d01817f76 --- /dev/null +++ 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 new file mode 100644 index 000000000000..265117bbc984 --- /dev/null +++ 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 new file mode 100644 index 000000000000..ebcc3f28c713 --- /dev/null +++ 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 new file mode 100644 index 000000000000..410f9edfa1ee --- /dev/null +++ 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 new file mode 100644 index 000000000000..8c0d49f735e1 --- /dev/null +++ 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 new file mode 100644 index 000000000000..9baf40c170f6 --- /dev/null +++ b/clang/examples/DPCT/sugar/kernel.cu @@ -0,0 +1,8 @@ +__global__ void f() {} + +void test() { + dim3 gridDim, blockDim; + // Start + f<<>>(); + // End +} diff --git a/clang/lib/DPCT/APIMapping/QueryAPIMapping.cpp b/clang/lib/DPCT/APIMapping/QueryAPIMapping.cpp index a5a40fd84cef..1527a2e0dacb 100644 --- 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 @@ -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 new file mode 100644 index 000000000000..c6515ecfd3b9 --- /dev/null +++ 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<<>>(); +// 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 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 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 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 index e8ae3e359981..f39767d7c319 100644 --- 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