From d536d58012923b8389ed745a8533846f767a98d6 Mon Sep 17 00:00:00 2001
From: intwanghao <hao3.wang@intel.com>
Date: Wed, 7 May 2025 13:29:45 +0800
Subject: [PATCH 1/2] fix

Signed-off-by: intwanghao <hao3.wang@intel.com>
---
 clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp     |  6 ++++
 .../RulesLang/RulesLangNoneAPIAndType.cpp     |  4 ++-
 clang/test/dpct/wmma2.cu                      | 35 +++++++++++++++++++
 3 files changed, 44 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/dpct/wmma2.cu

diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp
index 41aa106e5060..cddd7f9a3f45 100644
--- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp
+++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp
@@ -525,6 +525,12 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) {
               dyn_cast<NamespaceDecl>(Qualifier->getAsNamespace())) {
         CTSName = getNameSpace(NSD) + "::" + DRE->getNameInfo().getAsString();
       }
+    } else if (auto NA = Qualifier->getAsNamespaceAlias()) {
+      auto ND = NA->getNamespace();
+      if (ND && (ND->getName() == "wmma") &&
+          dpct::DpctGlobalInfo::isInCudaPath(ND->getBeginLoc())) {
+        CTSName = getNameSpace(ND) + "::" + DRE->getNameInfo().getAsString();
+      }
     } else if (!IsNamespaceOrAlias || !IsSpecicalAPI) {
       if (DRE->getDecl()->isCXXClassMember()) {
         std::string Result;
diff --git a/clang/lib/DPCT/RulesLang/RulesLangNoneAPIAndType.cpp b/clang/lib/DPCT/RulesLang/RulesLangNoneAPIAndType.cpp
index 027c92263670..bcc2e0c5e117 100644
--- a/clang/lib/DPCT/RulesLang/RulesLangNoneAPIAndType.cpp
+++ b/clang/lib/DPCT/RulesLang/RulesLangNoneAPIAndType.cpp
@@ -1123,8 +1123,10 @@ void NamespaceRule::runRule(const MatchFinder::MatchResult &Result) {
   } else if (auto NAD = getAssistNodeAsType<NamespaceAliasDecl>(
                  Result, "namespaceAlias")) {
     std::string Namespace = NAD->getNamespace()->getNameAsString();
-    if (Namespace == "cooperative_groups" || Namespace == "placeholders")
+    if (Namespace == "cooperative_groups" || Namespace == "placeholders" ||
+        Namespace == "wmma") {
       emplaceTransformation(new ReplaceDecl(NAD, ""));
+    }
   } else if (auto UD = getAssistNodeAsType<UsingDecl>(Result, "using")) {
     auto &SM = DpctGlobalInfo::getSourceManager();
     SourceLocation Beg, End;
diff --git a/clang/test/dpct/wmma2.cu b/clang/test/dpct/wmma2.cu
new file mode 100644
index 000000000000..eff6654e6057
--- /dev/null
+++ b/clang/test/dpct/wmma2.cu
@@ -0,0 +1,35 @@
+// clang-format off
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0
+// RUN: dpct --format-range=none --use-experimental-features=matrix -out-root %T/wmma2 %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/wmma2/wmma2.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST  %T/wmma2/wmma2.dp.cpp -o %T/wmma2/wmma2.dp.o %}
+
+#include <assert.h>
+#include <cuda.h>
+#include <iostream>
+#include <mma.h>
+// CHECK: #include <sycl/sycl.hpp>
+// CHECK: #include <dpct/dpct.hpp>
+
+// CHECK-NOT: ^namespace wmmaa = nvcuda::wmma;
+
+namespace wmmaa = nvcuda::wmma;
+
+template<typename T>
+__global__ void simple_wmma_gemm(T *d) {
+  wmmaa::fragment<wmmaa::accumulator, 16, 16, 16, T> c_frag;
+// CHECK: sycl::ext::oneapi::experimental::matrix::joint_matrix_store(sycl::ext::oneapi::this_work_item::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);
+  wmmaa::store_matrix_sync(d, c_frag, 1, wmmaa::mem_row_major);
+}
+
+int main() {
+
+  simple_wmma_gemm<half><<<1, 1>>>(nullptr);
+
+  simple_wmma_gemm<float><<<1, 1>>>(nullptr);
+
+  return 0;
+}
+
+

From 3d63f1ba755f2adb023ab9e815a08cef46aee706 Mon Sep 17 00:00:00 2001
From: intwanghao <hao3.wang@intel.com>
Date: Wed, 7 May 2025 15:01:58 +0800
Subject: [PATCH 2/2] fix

Signed-off-by: intwanghao <hao3.wang@intel.com>
---
 clang/test/dpct/wmma2.cu | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/clang/test/dpct/wmma2.cu b/clang/test/dpct/wmma2.cu
index eff6654e6057..bbf7feecde7f 100644
--- a/clang/test/dpct/wmma2.cu
+++ b/clang/test/dpct/wmma2.cu
@@ -11,9 +11,6 @@
 #include <mma.h>
 // CHECK: #include <sycl/sycl.hpp>
 // CHECK: #include <dpct/dpct.hpp>
-
-// CHECK-NOT: ^namespace wmmaa = nvcuda::wmma;
-
 namespace wmmaa = nvcuda::wmma;
 
 template<typename T>