revert custom kernels tiling on device action

Signed-off-by: leo-pony <nengjunma@outlook.com>

Author: leo-pony

.github/workflows/vllm_ascend_test.yaml

@@ -137,7 +137,7 @@ jobs:
       max-parallel: 2
       matrix:
         os: [linux-arm64-npu-1]
-        vllm_version: [v0.9.2]
+        vllm_version: [main, v0.9.2]
     name: singlecard e2e test
     runs-on: ${{ matrix.os }}
     container:
@@ -211,13 +211,13 @@ jobs:
           VLLM_USE_MODELSCOPE=True pytest -sv tests/e2e/singlecard/spec_decode_v1/test_v1_spec_decode.py
 
   e2e-4-cards:
-    needs: [lint]
-    if: ${{ needs.lint.result == 'success' }}
+    needs: [e2e]
+    if: ${{ needs.e2e.result == 'success' }}
     strategy:
       max-parallel: 1
       matrix:
         os: [linux-arm64-npu-4]
-        vllm_version: [v0.9.2]
+        vllm_version: [main, v0.9.2]
     name: multicard e2e test
     runs-on: ${{ matrix.os }}
     container:

README.md

@@ -42,7 +42,7 @@ By using vLLM Ascend plugin, popular open-source models, including Transformer-l
 - Software:
   * Python >= 3.9, < 3.12
   * CANN >= 8.1.RC1
-  * PyTorch >= 2.5.1, torch-npu >= 2.7.1rc1
+  * PyTorch >= 2.7.1, torch-npu >= 2.7.1rc1
   * vLLM (the same version as vllm-ascend)
 
 ## Getting Started

csrc/kernels/get_masked_input_and_mask_kernel.cpp

@@ -325,12 +325,13 @@ extern "C" __global__ __aicore__ void get_masked_input_and_mask_kernel(
     const int64_t num_org_vocab_padding,
     const int64_t added_vocab_start_index,
     const int64_t added_vocab_end_index,
-    const int64_t size)
+    const int64_t size,
+    const uint32_t loop_cnt,
+    const uint32_t aiv_num)
 {
     {
         GetMaskedInputAndMask<int32_t> op{};
-        uint32_t aiv_num = AscendC::GetBlockNum();
-        uint32_t loop_cnt = (size + aiv_num - 1) / aiv_num;
+
         for (int64_t i = AscendC::GetBlockIdx(); i < loop_cnt; i += aiv_num) {
             op.Init(input + i * size/loop_cnt, 
                    masked_input + i * size/loop_cnt,
@@ -356,11 +357,11 @@ void get_masked_input_and_mask_impl(
     const int64_t num_org_vocab_padding, 
     const int64_t added_vocab_start_index,
     const int64_t added_vocab_end_index,
-    const int64_t size)
+    const int64_t size,
+    const uint32_t loop_cnt,
+    const uint32_t aiv_num)
 {
-    // block_dim only used for parameter check, execute parallel number is handled in get_masked_input_and_mask_kernel
-    const uint32_t block_dim = 8;
-    get_masked_input_and_mask_kernel<<<block_dim, nullptr, stream>>>(
+    get_masked_input_and_mask_kernel<<<aiv_num, nullptr, stream>>>(
         static_cast<int32_t*>(input),
         static_cast<int32_t*>(masked_input),
         static_cast<bool*>(mask_out),
@@ -369,7 +370,9 @@ void get_masked_input_and_mask_impl(
         num_org_vocab_padding,
         added_vocab_start_index,
         added_vocab_end_index,
-        size);
+        size,
+        loop_cnt,
+        aiv_num);
 }
 
 } // namespace vllm_ascend

csrc/kernels/pos_encoding_kernels.cpp

@@ -311,7 +311,7 @@ template <typename scalar_t, bool isNeox> class RotaryEmbedding {
         __gm__ int64_t* positions, __gm__ void* queryDst, __gm__ void* keyDst, __gm__ TYPE* query, __gm__ TYPE* key,            \
         __gm__ TYPE* cosSinCache, const int rotDim, const int64_t queryStride, const int64_t keyStride,                         \
         const int64_t dstQueryStride, const int64_t dstKeyStride, const int numHeads, const int numKvHeads,                     \
-        const int headSize, const int64_t numTokens, const int coreNum)                                      \
+        const int headSize, const int64_t numTokens, const int loopNum, const int coreNum)                                      \
     {                                                                                                                           \
         AscendC::TPipe pipe;                                                                                                    \
         RotaryEmbedding<TYPE, NEOX> op{};                                                                                       \
@@ -341,12 +341,12 @@ namespace vllm_ascend {
         rope_custom_true_##TYPE<<<blockDim, nullptr, stream>>>(                                                  \
             positions, queryDst, keyDst, reinterpret_cast<TYPE *>(query), reinterpret_cast<TYPE *>(key),         \
             reinterpret_cast<TYPE *>(cosSinCache), rotDim, queryStride, keyStride, dstQueryStride, dstKeyStride, \
-            numHeads, numKvHeads, headSize, numTokens, blockDim);                                       \
+            numHeads, numKvHeads, headSize, numTokens, loopCnt, blockDim);                                       \
     else                                                                                                         \
         rope_custom_false_##TYPE<<<blockDim, nullptr, stream>>>(                                                 \
             positions, queryDst, keyDst, reinterpret_cast<TYPE *>(query), reinterpret_cast<TYPE *>(key),         \
             reinterpret_cast<TYPE *>(cosSinCache), rotDim, queryStride, keyStride, dstQueryStride, dstKeyStride, \
-            numHeads, numKvHeads, headSize, numTokens, blockDim);
+            numHeads, numKvHeads, headSize, numTokens, loopCnt, blockDim);
 
 // maximum number for runtime to launch a ascendc kernel.
 // we use this to constrain the maximum number of block size
@@ -356,7 +356,8 @@ extern void rotary_embedding_impl(AscendType type, bool isNeox, void *stream, in
                                     void *keyDst, void *query, void *key, void *cosSinCache, const int rotDim,
                                     const int64_t queryStride, const int64_t keyStride, const int64_t dstQueryStride,
                                     const int64_t dstKeyStride, const int numHeads, const int numKvHeads,
-                                    const int headSize, const int64_t numTokens)
+                                    const int headSize, const int64_t numTokens, const uint32_t loopCnt,
+                                    uint32_t aivNum)
 {
 
     int blockDim = maxParallelSize > numTokens ? numTokens : maxParallelSize;

csrc/ops.h

@@ -28,7 +28,8 @@ namespace vllm_ascend {
     void *keyDst, void *query, void *key, void *cosSinCache, const int rotDim,
     const int64_t queryStride, const int64_t keyStride, const int64_t dstQueryStride,
     const int64_t dstKeyStride, const int numHeads, const int numKvHeads,
-    const int headSize, const int64_t numTokens);
+    const int headSize, const int64_t numTokens, const uint32_t loopCnt,
+    uint32_t aivNum);
 
   extern void get_masked_input_and_mask_impl(
     void* stream,
@@ -40,7 +41,9 @@ namespace vllm_ascend {
     const int64_t num_org_vocab_padding, 
     const int64_t added_vocab_start_index,
     const int64_t added_vocab_end_index,
-    const int64_t size);
+    const int64_t size,
+    const uint32_t loop_cnt,
+    const uint32_t aiv_num);
 
   torch::Tensor weak_ref_tensor(torch::Tensor& tensor) {
     if (!tensor.is_privateuseone()) {

docs/source/installation.md

@@ -13,7 +13,7 @@ This document describes how to install vllm-ascend manually.
     |---------------|----------------------------------|-------------------------------------------|
     | CANN          | >= 8.1.RC1                       | Required for vllm-ascend and torch-npu    |
     | torch-npu     | >= 2.7.1rc1                      | Required for vllm-ascend, No need to install manually, it will be auto installed in below steps |
-    | torch         | >= 2.5.1                         | Required for torch-npu and vllm           |
+    | torch         | >= 2.7.1                         | Required for torch-npu and vllm           |
 
 You have 2 way to install:
 - **Using pip**: first prepare env manually or via CANN image, then install `vllm-ascend` using pip.