[SYCLomatic] Fix a bug that wmma APIs not migrated in template function (#2424)

Signed-off-by: Ziran Zhang <[email protected]>

Author: ziranzha

clang/lib/DPCT/ExprAnalysis.cpp

@@ -859,15 +859,7 @@ void ExprAnalysis::analyzeExpr(const UnaryExprOrTypeTraitExpr *UETT) {
 }
 
 inline void ExprAnalysis::analyzeExpr(const UnresolvedLookupExpr *ULE) {
-  RefString.clear();
-  llvm::raw_string_ostream OS(RefString);
-  if (auto NNS = ULE->getQualifier()) {
-    if (NNS->getKind() != clang::NestedNameSpecifier::SpecifierKind::Global) {
-      NNS->print(OS, dpct::DpctGlobalInfo::getContext().getPrintingPolicy());
-    }
-  }
-  ULE->getName().print(OS,
-                       dpct::DpctGlobalInfo::getContext().getPrintingPolicy());
+  RefString = ULE->decls().begin().getDecl()->getQualifiedNameAsString();
 }
 
 void ExprAnalysis::analyzeExpr(const ExplicitCastExpr *Cast) {

clang/lib/DPCT/WMMAAPIMigration.cpp

@@ -13,17 +13,22 @@ using namespace clang::dpct;
 using namespace clang::ast_matchers;
 
 void clang::dpct::WMMARule::registerMatcher(ast_matchers::MatchFinder &MF) {
+  auto FuncName = []() {
+    return hasAnyName("fill_fragment", "load_matrix_sync", "mma_sync",
+                      "store_matrix_sync");
+  };
   MF.addMatcher(
       typeLoc(loc(qualType(hasDeclaration(namedDecl(allOf(
                   hasAnyName("fragment", "matrix_a", "matrix_b", "row_major",
                              "col_major", "accumulator", "layout_t"),
                   hasDeclContext(namespaceDecl(hasName("wmma")))))))))
           .bind("type"),
       this);
-  MF.addMatcher(callExpr(callee(functionDecl(allOf(
-                             hasAnyName("fill_fragment", "load_matrix_sync",
-                                        "mma_sync", "store_matrix_sync"),
-                             hasDeclContext(namespaceDecl(hasName("wmma")))))))
+  MF.addMatcher(callExpr(anyOf(callee(functionDecl(allOf(
+                                   FuncName(), hasDeclContext(namespaceDecl(
+                                                   hasName("wmma")))))),
+                               callee(unresolvedLookupExpr(
+                                   hasAnyDeclaration(namedDecl(FuncName()))))))
                     .bind("call"),
                 this);
   MF.addMatcher(declRefExpr(to(enumConstantDecl(allOf(

clang/test/dpct/wmma.cu

@@ -201,5 +201,19 @@ int main() {
 
   return 0;
 }
+
+using namespace nvcuda;
+template<typename T>
+__global__ void simple_wmma_gemm(T *d) {
+  wmma::fragment<wmma::accumulator, 16, 16, 16, T> c_frag;
+  // CHECK: sycl::ext::oneapi::experimental::matrix::joint_matrix_store(item_ct1.get_sub_group(), c_frag.get(), sycl::address_space_cast<sycl::access::address_space::generic_space, sycl::access::decorated::no, T>(d), 1, sycl::ext::oneapi::experimental::matrix::layout::row_major);
+  wmma::store_matrix_sync(d, c_frag, 1, wmma::mem_row_major);
+}
+int main() {
+  simple_wmma_gemm<half><<<1, 1>>>(nullptr);
+  simple_wmma_gemm<float><<<1, 1>>>(nullptr);
+  return 0;
+}
+
 // clang-format on
 #endif