[BugFix]This PR aims to fix the precision issue of the LoRA feature in vllm-ascend.

Co-authored-by: liuchenbing <[email protected]>
Co-authored-by: guanyuzhu <[email protected]>
vLLM version: v0.11.0
vLLM main: vllm-project/vllm
signed-off-by: hukongyi <[email protected]>

Author: 3 people

csrc/kernels/bgmv_expand.cpp

@@ -342,7 +342,7 @@ class BGMVExpand {
 
 // declare all dtype kernel
 BGMV_EXPAND_TYPE_DECLARE(half)
-#if (__CCE_AICORE__ >= 220)
+#if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
     BGMV_EXPAND_TYPE_DECLARE(bfloat16_t)
 #endif
 
@@ -356,9 +356,11 @@ extern void bgmv_expand_impl(AscendType type, void* stream, void* x, void* weigh
         bgmv_expand_half<<<blockDim, nullptr, stream>>>(x, weight, indices, indicesSize, yIn, yOut, batchSize, numTokensPerCore,
                                                         maxLoRARank, outputHiddenDim, sliceOffset, outputFullDim);
     } else if (type == AscendType::BF16) {
+        #if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
         bgmv_expand_bfloat16_t<<<blockDim, nullptr, stream>>>(x, weight, indices, indicesSize, yIn, yOut, batchSize,
                                                                   numTokensPerCore, maxLoRARank, outputHiddenDim,
                                                                   sliceOffset, outputFullDim);
+        #endif
     } else {
         return;
     }

csrc/kernels/bgmv_shrink.cpp

@@ -226,7 +226,7 @@ class BGMVShrink {
 
 // declare all dtype kernel
 BGMV_SHRINK_TYPE_DECLARE(half)
-#if (__CCE_AICORE__ >= 220)
+#if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
     BGMV_SHRINK_TYPE_DECLARE(bfloat16_t)
 #endif
 
@@ -240,8 +240,10 @@ extern void bgmv_shrink_impl(AscendType type, void* stream, void* x, void* weigh
         bgmv_shrink_half<<<blockDim, nullptr, stream>>>(x, weight, indices, indicesSize, y, batchSize, numTokensPerCore, 
                                                         inputHiddenDim, maxLoRARank, scale);
     } else if (type == AscendType::BF16) {
+        #if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
         bgmv_shrink_bfloat16_t<<<blockDim, nullptr, stream>>>(x, weight, indices, indicesSize, y, batchSize, numTokensPerCore, 
                                                                   inputHiddenDim, maxLoRARank, scale);
+        #endif
     } else {
         return;
     }

csrc/kernels/sgmv_expand.cpp

@@ -357,7 +357,7 @@ class SGMVExpand {
 
 // declare all dtype kernel
 SGMV_EXPAND_TYPE_DECLARE(half)
-#if (__CCE_AICORE__ >= 220)
+#if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
     SGMV_EXPAND_TYPE_DECLARE(bfloat16_t)
 #endif
 
@@ -375,10 +375,12 @@ extern void sgmv_expand_impl(AscendType type, void* stream, void* x, void* weigh
                                                         numTokensPerCore, maxLoRARank, outputHiddenDim, sliceOffset, 
                                                         outputFullDim);
     } else if (type == AscendType::BF16) {
+        #if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
         sgmv_expand_bfloat16_t<<<blockDim, nullptr, stream>>>(x, weight, loraIndices, loraIndicesSize, 
                                                                   seqLen, seqLenSize, yIn, yOut, batchSize,
                                                                   numTokensPerCore, maxLoRARank, outputHiddenDim,
                                                                   sliceOffset, outputFullDim);
+        #endif
     } else {
         return;
     }

csrc/kernels/sgmv_shrink.cpp

@@ -242,7 +242,7 @@ class SGMVShrink {
 
 // declare all dtype kernel
 SGMV_SHRINK_TYPE_DECLARE(half)
-#if (__CCE_AICORE__ >= 220)
+#if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
     SGMV_SHRINK_TYPE_DECLARE(bfloat16_t)
 #endif
 
@@ -260,11 +260,13 @@ extern void sgmv_shrink_impl(AscendType type, void* stream, void* x, void* weigh
                                                         numTokensPerCore, inputHiddenDim, maxLoRARank,
                                                         scale);
     } else if (type == AscendType::BF16) {
+        #if !defined(__CCE_AICORE__) || (__CCE_AICORE__ >= 220)
         sgmv_shrink_bfloat16_t<<<blockDim, nullptr, stream>>>(x, weight, loraIndices, loraIndicesSize, 
                                                                   seqLen, seqLenSize, 
                                                                   y, batchSize,
                                                                   numTokensPerCore, inputHiddenDim, maxLoRARank,
                                                                   scale);
+        #endif
     } else {
         return;
     }