From e6ee3c62c46eae7cd9b01d992548444e8dcb5213 Mon Sep 17 00:00:00 2001 From: Yihan Wang Date: Wed, 26 Jun 2024 10:12:32 +0800 Subject: [PATCH] [SYCLomatic] Support migration of cub::{SmVersion, SmVersionUncached} (#2091) Signed-off-by: Wang, Yihan --- clang/lib/DPCT/APINames_CUB.inc | 4 +-- clang/lib/DPCT/CUBAPIMigration.cpp | 3 ++- .../CUB/RewriterUtilityFunctions.cpp | 26 ++++++++++++++++++- clang/test/dpct/cub/intrinsic/sm_version.cu | 19 ++++++++++++++ 4 files changed, 48 insertions(+), 4 deletions(-) create mode 100644 clang/test/dpct/cub/intrinsic/sm_version.cu diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc index 8e45f4c72d21..b6e975215f18 100644 --- 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 index c9f91d9ae37b..d32494711f61 100644 --- 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 index 241f39869ab7..f1a89bf9cd2b 100644 --- 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 new file mode 100644 index 000000000000..261c15731373 --- /dev/null +++ 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 + +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); +}