From 47f0b52fe592a2185c626e1d97b06e5573665820 Mon Sep 17 00:00:00 2001
From: Li-Jian <lijian120@huawei.com>
Date: Thu, 11 Sep 2025 11:39:44 +0800
Subject: [PATCH] fix precision problem on david

---
 .../MatmulInvocationNeo/main.cpp                  |  4 ++--
 .../MatmulInvocationNeo/matmul_custom.cpp         |  3 +++
 .../MatmulInvocationNeo/matmul_custom_tiling.cpp  | 15 +++++++++------
 3 files changed, 14 insertions(+), 8 deletions(-)

diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp
index 17518d258..1fb5a80d8 100644
--- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp
+++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/main.cpp
@@ -34,9 +34,9 @@ int32_t main(int32_t argc, char *argv[])
     uint8_t *tilingBuf = (uint8_t *)malloc(tilingFileSize);
     GenerateTiling(socVersion, tilingBuf);
 #ifdef CUSTOM_ASCEND310P
-    uint32_t blockDim = 2;
+    uint32_t blockDim = reinterpret_cast<TCubeTiling *>(tilingBuf)->usedCoreNum;
 #else
-    uint32_t blockDim = 22;
+    uint32_t blockDim = (reinterpret_cast<TCubeTiling *>(tilingBuf)->usedCoreNum + 1) / 2;
 #endif
 
 #ifdef ASCENDC_CPU_DEBUG
diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom.cpp b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom.cpp
index 19292d6c4..f3a763b8d 100644
--- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom.cpp
+++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom.cpp
@@ -120,6 +120,9 @@ extern "C" __global__ __aicore__ void matmul_custom(GM_ADDR a, GM_ADDR b, GM_ADD
            MatmulType<AscendC::TPosition::GM, CubeFormat::ND, B_T>,
            MatmulType<AscendC::TPosition::GM, CubeFormat::ND, C_T>> mm;
     REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm, &tiling); // Initialize the matmul object.
+    if (GetBlockIdx() >= tiling.usedCoreNum) {
+        return;
+    }
 #ifdef CUSTOM_ASCEND310P
     // Set temp UB space when on ASCEND310P.
     AscendC::TBuf<> tmpMMFormatUb;
diff --git a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom_tiling.cpp b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom_tiling.cpp
index d3f0dc2db..f86f9b3c0 100644
--- a/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom_tiling.cpp
+++ b/operator/ascendc/0_introduction/11_matmul_kernellaunch/MatmulInvocationNeo/matmul_custom_tiling.cpp
@@ -25,9 +25,9 @@ using namespace std;
   */
 void GenerateTiling(const char *socVersion, uint8_t *tilingBuf)
 {
-    int M = 512;
-    int N = 1024;
-    int K = 512;
+    constexpr int32_t M = 512;
+    constexpr int32_t N = 1024;
+    constexpr int32_t K = 512;
 
     TPosition leftPosition = TPosition::GM;
     CubeFormat leftFormat = CubeFormat::ND;
@@ -45,8 +45,8 @@ void GenerateTiling(const char *socVersion, uint8_t *tilingBuf)
 
     bool isBias = false;
 
-    int32_t singleCoreM = 512;
-    int32_t singleCoreN = 512;
+    constexpr int32_t SINGLECORE_M = 512;
+    constexpr int32_t SINGLECORE_N = 512;
 
     optiling::TCubeTiling tilingData;
     auto ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance(socVersion);
@@ -60,7 +60,10 @@ void GenerateTiling(const char *socVersion, uint8_t *tilingBuf)
     tilingApi.SetOrgShape(M, N, K);
     tilingApi.SetShape(M, N, K);
     if (ascendcPlatform->GetSocVersion() == platform_ascendc::SocVersion::ASCEND310P) {
-        tilingApi.SetSingleShape(singleCoreM, singleCoreN, -1);  // Set the fixed singleCoreM=512, singleCoreN=512.
+        tilingApi.SetSingleShape(SINGLECORE_M, SINGLECORE_N, -1);  // Set the fixed singleCoreM=512, singleCoreN=512.
+        int32_t mBlockNum = M / SINGLECORE_M;
+        int32_t nBlockNum = N / SINGLECORE_N;
+        tilingApi.SetDim(mBlockNum * nBlockNum);
     }
     tilingApi.SetBias(isBias);
     tilingApi.SetBufferSpace(-1, -1, -1);
-- 
Gitee