[SYCLomatic] Support migration of __cvta_generic_to_shared

clang/lib/DPCT/ASTTraversal.cpp

@@ -155,6 +155,7 @@ REGISTER_RULE(CompatWithClangRule, PassKind::PK_Migration)
 REGISTER_RULE(AssertRule, PassKind::PK_Migration)
 REGISTER_RULE(GraphRule, PassKind::PK_Migration)
 REGISTER_RULE(GraphicsInteropRule, PassKind::PK_Migration)
+REGISTER_RULE(RulesLangAddrSpaceConvRule, PassKind::PK_Migration)
 
 REGISTER_RULE(BLASEnumsRule, PassKind::PK_Migration, RuleGroupKind::RK_BLas)
 REGISTER_RULE(BLASFunctionCallRule, PassKind::PK_Migration,RuleGroupKind::RK_BLas)

clang/lib/DPCT/CMakeLists.txt

@@ -119,6 +119,7 @@ add_clang_library(DPCT
   RulesLang/RulesLangAtomic.cpp
   RulesLang/RulesLangCooperativeGroups.cpp
   RulesLang/RulesLangTexture.cpp
+  RulesLang/RulesLangAddrSpaceConv.cpp
   RulesLangLib/CUB/CallExprRewriterCUB.cpp
   RulesLangLib/CUB/RewriterClassMethods.cpp
   RulesLangLib/CUB/RewriterDeviceHistgram.cpp

clang/lib/DPCT/RulesLang/RulesLang.h

@@ -997,6 +997,13 @@ class GraphicsInteropRule : public NamedMigrationRule<GraphicsInteropRule> {
   void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
 };
 
+class RulesLangAddrSpaceConvRule
+    : public NamedMigrationRule<RulesLangAddrSpaceConvRule> {
+public:
+  void registerMatcher(ast_matchers::MatchFinder &MF) override;
+  void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
+};
+
 TextModification *replaceText(SourceLocation Begin, SourceLocation End,
                               std::string &&Str, const SourceManager &SM);
 

clang/lib/DPCT/RulesLang/RulesLangAddrSpaceConv.cpp

@@ -0,0 +1,77 @@
+//===--------------- RulesLangAddrSpaceConv.cpp ---------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "AnalysisInfo.h"
+#include "RuleInfra/ExprAnalysis.h"
+#include "RulesLang.h"
+#include "Utility.h"
+#include "clang/AST/Decl.h"
+#include "clang/AST/Expr.h"
+#include "clang/AST/Stmt.h"
+#include "clang/ASTMatchers/ASTMatchers.h"
+#include "llvm/Support/Casting.h"
+
+namespace clang {
+namespace dpct {
+
+using namespace clang::ast_matchers;
+
+void RulesLangAddrSpaceConvRule::registerMatcher(MatchFinder &MF) {
+  MF.addMatcher(
+      callExpr(callee(functionDecl(hasName("__cvta_generic_to_shared"))))
+          .bind("call"),
+      this);
+}
+
+void RulesLangAddrSpaceConvRule::runRule(
+    const MatchFinder::MatchResult &Result) {
+  const auto *CE = getNodeAsType<CallExpr>(Result, "call");
+  if (!CE)
+    return;
+  // Check if meets below conditions:
+  // (1) A vardecl's init value is the "call" (or after type cast).
+  // (2) The var is only used as the asm stmt parameter.
+
+  // Check (1)
+  const auto *DS = DpctGlobalInfo::findAncestor<DeclStmt>(CE);
+  if (!DS)
+    return;
+  const auto *VD =
+      DS->isSingleDecl() ? dyn_cast<VarDecl>(DS->getSingleDecl()) : nullptr;
+  if (!VD)
+    return;
+  const auto *Init = VD->getInit();
+  if (!Init)
+    return;
+  if (Init->IgnoreCasts() != CE)
+    return;
+
+  // Check (2)
+  const auto *Ctx = VD->getDeclContext();
+  const auto *FD = dyn_cast<FunctionDecl>(Ctx);
+  if (!FD)
+    return;
+  if (!FD->hasBody())
+    return;
+  std::set<const clang::DeclRefExpr *> DREs =
+      matchTargetDREInScope(VD, FD->getBody());
+  if (DREs.size() != 1)
+    return;
+  const auto *DRE = *DREs.begin();
+  const auto *AS = DpctGlobalInfo::findAncestor<AsmStmt>(DRE);
+  if (!AS)
+    return;
+
+  // Generate replacement
+  std::string ReplacementStr = "auto " + VD->getNameAsString() + " = " +
+                               ExprAnalysis::ref(CE->getArg(0)) + ";";
+  emplaceTransformation(new ReplaceDecl(VD, std::move(ReplacementStr)));
+}
+
+} // namespace dpct
+} // namespace clang

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -2308,6 +2308,16 @@ ENTRY(make_cuFloatComplex, make_cuFloatComplex, true, NO_FLAG, P4, "comment")
 ENTRY(__assert_fail, __assert_fail, true, NO_FLAG, P4, "Successful")
 ENTRY(__assertfail, __assertfail, true, NO_FLAG, P4, "Successful")
 
+// Address Space Conversion Functions
+ENTRY(__cvta_generic_to_global, __cvta_generic_to_global, false, NO_FLAG, P7, "comment")
+ENTRY(__cvta_generic_to_shared, __cvta_generic_to_shared, true, NO_FLAG, P7, "comment")
+ENTRY(__cvta_generic_to_constant, __cvta_generic_to_constant, false, NO_FLAG, P7, "comment")
+ENTRY(__cvta_generic_to_local, __cvta_generic_to_local, false, NO_FLAG, P7, "comment")
+ENTRY(__cvta_global_to_generic, __cvta_global_to_generic, false, NO_FLAG, P7, "comment")
+ENTRY(__cvta_shared_to_generic, __cvta_shared_to_generic, false, NO_FLAG, P7, "comment")
+ENTRY(__cvta_constant_to_generic, __cvta_constant_to_generic, false, NO_FLAG, P7, "comment")
+ENTRY(__cvta_local_to_generic, __cvta_local_to_generic, false, NO_FLAG, P7, "comment")
+
 ENTRY(cuGetExportTable, cuGetExportTable, true, NO_FLAG, P7, "Partial")
 
 // clang-format on

clang/test/dpct/addr_space_conv.cu

@@ -0,0 +1,25 @@
+// 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 --out-root %T/addr_space_conv %s --cuda-include-path="%cuda-path/include"
+// RUN: FileCheck --input-file %T/addr_space_conv/addr_space_conv.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -DNO_BUILD_TEST -fsycl %T/addr_space_conv/addr_space_conv.dp.cpp -o %T/addr_space_conv/addr_space_conv.dp.o %}
+
+#include <cstdint>
+
+__global__ void kernel1(const void* ptr) {
+  // In PTX, addresses of the local and shared memory spaces are always 32 bits in size.
+  __shared__ float shared_array[1024];
+  // CHECK: auto smem = shared_array;
+  uint32_t smem = static_cast<uint32_t>(__cvta_generic_to_shared(shared_array));
+#ifndef NO_BUILD_TEST
+  asm volatile(
+    "{\n"
+    "   cp.async.cg.shared.global [%0], [%1], 16;\n"
+    "}\n" :: "r"(smem), "l"(ptr)
+  );
+#endif
+}
+
+void foo(const void* ptr) {
+  kernel1<<<1, 1>>>(ptr);
+}