[SYCLomatic] Migrate the LibCU APIs in the template class. (#2880)

Signed-off-by: Chen, Sheng S <sheng.s.chen@intel.com>

Author: ShengchenJ

clang/lib/DPCT/RulesLangLib/LIBCUAPIMigration.cpp

@@ -37,6 +37,13 @@ void LIBCURule::registerMatcher(ast_matchers::MatchFinder &MF) {
                         "compare_exchange_strong", "fetch_add", "fetch_sub",
                         "at");
     };
+    auto LIBCUMemberHasNames = [&]() {
+      return anyOf(
+          hasMemberName("load"), hasMemberName("store"),
+          hasMemberName("exchange"), hasMemberName("compare_exchange_weak"),
+          hasMemberName("compare_exchange_strong"), hasMemberName("fetch_add"),
+          hasMemberName("fetch_sub"), hasMemberName("at"));
+    };
     auto LIBCUTypesHasNames = [&]() {
       return hasAnyName("cuda::atomic", "cuda::std::atomic",
                         "cuda::std::array");
@@ -47,6 +54,9 @@ void LIBCURule::registerMatcher(ast_matchers::MatchFinder &MF) {
                             callee(cxxMethodDecl(LIBCUMemberFuncHasNames()))))
                       .bind("MemberCall"),
                   this);
+    MF.addMatcher(cxxDependentScopeMemberExpr(LIBCUMemberHasNames())
+                      .bind("DependentMemCall"),
+                  this);
   }
   {
     MF.addMatcher(dependentScopeDeclRefExpr().bind("DependentScope"),
@@ -88,6 +98,16 @@ void LIBCURule::runRule(const ast_matchers::MatchFinder::MatchResult &Result) {
   if (const CXXMemberCallExpr *MC =
           getNodeAsType<CXXMemberCallExpr>(Result, "MemberCall")) {
     EA.analyze(MC);
+  } else if (const CXXDependentScopeMemberExpr *CDSE =
+                 getNodeAsType<CXXDependentScopeMemberExpr>(
+                     Result, "DependentMemCall")) {
+    auto Parent = dpct::DpctGlobalInfo::getContext().getParents(*CDSE);
+    auto *CE = Parent[0].get<CallExpr>();
+    if (CE) {
+      for (size_t i = 0; i < CE->getNumArgs(); i++) {
+        EA.analyze(CE->getArg(i));
+      }
+    }
   } else if (const CallExpr *CE = getNodeAsType<CallExpr>(Result, "FuncCall")) {
     EA.analyze(CE);
   } else if (auto TL = getNodeAsType<TypeLoc>(Result, "TypeLoc")) {

clang/test/dpct/LibCU/libcu_atomic.cu

@@ -10,6 +10,24 @@
 // CHECK: #include <dpct/atomic.hpp>
 #include <cuda/atomic>
 
+template <class T> bool is_complete(const T &result) {
+  return !(result == static_cast<T>(0.0) && std::signbit(result));
+}
+
+template <typename T> struct ReduceArg {
+
+private:
+// CHECK:  dpct::atomic<T, sycl::memory_scope::system, sycl::memory_order::relaxed> *result_h;
+  cuda::atomic<T, cuda::thread_scope_system> *result_h;
+
+public:
+  void complete() {
+    //CHECK: while (!is_complete(result_h[0].load(sycl::memory_order::relaxed))) {
+    while (!is_complete(result_h[0].load(cuda::std::memory_order_relaxed))) {
+    }
+  }
+};
+
 int main(){
   // CHECK: sycl::atomic_fence(sycl::memory_order::release, sycl::memory_scope::system);
   cuda::atomic_thread_fence(cuda::std::memory_order_release, cuda::thread_scope_system);