From da4a124114a583fd7986f9d0dc9aee89f8cbed00 Mon Sep 17 00:00:00 2001 From: sunshine Date: Thu, 30 Oct 2025 15:58:30 +0800 Subject: [PATCH] Tiling Key Adapted to AscendType Enum --- kernels/fused_rope/fused_rope.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/kernels/fused_rope/fused_rope.cpp b/kernels/fused_rope/fused_rope.cpp index 63396e1..66e105e 100644 --- a/kernels/fused_rope/fused_rope.cpp +++ b/kernels/fused_rope/fused_rope.cpp @@ -1,6 +1,7 @@ #include "kernel_operator.h" #include "fused_rope_bf16.h" #include "fused_rope_fp32.h" +#include "../types.h" using namespace AscendC; using namespace FusedRope; @@ -18,7 +19,7 @@ extern "C" __global__ __aicore__ void FusedRopeKernel( { TPipe pipe; // DT_BF16 - if (tilingKey == 20) { + if (tilingKey == (uint64_t)kvcache_ops::AscendType::BF16) { TPipe* ptr = &pipe; if (ptr != nullptr) { FusedRopeFP16 op; @@ -35,7 +36,7 @@ extern "C" __global__ __aicore__ void FusedRopeKernel( } } // DT_FLOAT16 - if (tilingKey == 21) { + if (tilingKey == (uint64_t)kvcache_ops::AscendType::FP16) { TPipe* ptr = &pipe; if (ptr != nullptr) { FusedRopeFP16 op; @@ -52,7 +53,7 @@ extern "C" __global__ __aicore__ void FusedRopeKernel( } } // DT_FLOAT - if (tilingKey == 22) { + if (tilingKey == (uint64_t)kvcache_ops::AscendType::FP32) { TPipe* ptr = &pipe; if (ptr != nullptr) { FusedRopeFP32 op; -- Gitee