[SYCLomatic] Enable lit case LibCU/libcu_atomic_using_ns.cu (#2230)

Signed-off-by: Jiang, Zhiwei <[email protected]>

Author: zhiweij1

clang/lib/DPCT/ExprAnalysis.cpp

@@ -26,6 +26,7 @@
 #include "clang/AST/StmtObjC.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/TypeLoc.h"
+#include "llvm/Support/raw_ostream.h"
 
 extern clang::tooling::UnifiedPath DpctInstallPath;
 namespace clang {
@@ -442,7 +443,7 @@ void ExprAnalysis::dispatch(const Stmt *Expression) {
   }
 }
 
-bool isMathFunction(std::string Name) {
+bool isMathFunctionExceptRewriter(std::string Name) {
   static std::set<std::string> MathFunctions = {
 #define ENTRY_RENAMED(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
 #define ENTRY_RENAMED_NO_REWRITE(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
@@ -467,6 +468,31 @@ bool isMathFunction(std::string Name) {
   return MathFunctions.count(Name);
 }
 
+bool isMathFunction(std::string Name) {
+  static std::set<std::string> MathFunctions = {
+#define ENTRY_RENAMED(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
+#define ENTRY_RENAMED_NO_REWRITE(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
+#define ENTRY_RENAMED_SINGLE(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
+#define ENTRY_RENAMED_DOUBLE(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
+#define ENTRY_EMULATED(SOURCEAPINAME, TARGETAPINAME) SOURCEAPINAME,
+#define ENTRY_OPERATOR(APINAME, OPKIND) APINAME,
+#define ENTRY_TYPECAST(APINAME) APINAME,
+#define ENTRY_UNSUPPORTED(APINAME) APINAME,
+#define ENTRY_REWRITE(APINAME) APINAME,
+#include "APINamesMath.inc"
+#undef ENTRY_RENAMED
+#undef ENTRY_RENAMED_NO_REWRITE
+#undef ENTRY_RENAMED_SINGLE
+#undef ENTRY_RENAMED_DOUBLE
+#undef ENTRY_EMULATED
+#undef ENTRY_OPERATOR
+#undef ENTRY_TYPECAST
+#undef ENTRY_UNSUPPORTED
+#undef ENTRY_REWRITE
+  };
+  return MathFunctions.count(Name);
+}
+
 bool isCGAPI(std::string Name) {
   return MapNames::CooperativeGroupsAPISet.count(Name);
 }
@@ -480,9 +506,10 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) {
             clang::NestedNameSpecifier::SpecifierKind::Namespace ||
         Qualifier->getKind() ==
             clang::NestedNameSpecifier::SpecifierKind::NamespaceAlias;
-    bool IsSpecicalAPI = isMathFunction(DRE->getNameInfo().getAsString()) ||
-                         isCGAPI(DRE->getNameInfo().getAsString());
-                         // for thrust::log10 and thrust::sinh ...
+    bool IsSpecicalAPI =
+        isMathFunctionExceptRewriter(DRE->getNameInfo().getAsString()) ||
+        isCGAPI(DRE->getNameInfo().getAsString());
+    // for thrust::log10 and thrust::sinh ...
     if (Qualifier->getKind() == NestedNameSpecifier::TypeSpec)
       analyzeType(DRE->getQualifierLoc().getTypeLoc());
     // log10 is a math function
@@ -513,6 +540,12 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) {
                   DRE->getNameInfo().getAsString();
       }
     }
+  } else if (const auto *FD = dyn_cast_or_null<FunctionDecl>(DRE->getDecl())) {
+    if (!isMathFunction(DRE->getNameInfo().getAsString()) &&
+        !isCGAPI(DRE->getNameInfo().getAsString())) {
+      llvm::raw_string_ostream OS(CTSName);
+      FD->printQualifiedName(OS);
+    }
   }
 
   if (!CTSName.empty()) {

clang/test/dpct/LibCU/libcu_atomic_using_ns.cu

@@ -1,5 +1,3 @@
-// UNSUPPORTED: system-linux
-// UNSUPPORTED: system-windows
 // UNSUPPORTED: v7.0, v7.5, v8.0, v9.0, v9.2, v10.0, v10.1, v10.2
 // UNSUPPORTED: cuda-7.0, cuda-7.5, cuda-8.0, cuda-9.0, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
 // RUN: dpct --format-range=none -in-root %S -out-root %T/Libcu %S/libcu_atomic_using_ns.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
@@ -11,8 +9,12 @@
 // CHECK: #include <dpct/atomic.hpp>
 #include <cuda/atomic>
 
-// CHECK-EMPTY
+// CHECK: // BEGIN
+// CHECK-EMPTY:
+// CHECK-NEXT: // END
+// BEGIN
 using namespace cuda;
+// END
 
 int main() {
   // CHECK: sycl::atomic_fence(sycl::memory_order::release, sycl::memory_scope::system);