8d75652949
Sources cloned and tree'd (no --depth):
- jd-opensource/xllm: ILU kernels, CUDA kernels, MoE kernels
- NVIDIA/cccl: CUB tuning/dispatch headers (block-level primitives)
- fla-org/flash-linear-attention: Triton GDN kernels
- NVIDIA/cutlass: grouped GEMM reference (read, not copied)
- Dao-AILab/flash-attention: attention kernel reference (SM80+, read only)
New CUDA kernels (from xllm, SM-agnostic, portable to BI-V100):
ex_engine/xllm_kernels/cuda/activation.cu (188 lines) — silu_and_mul, gelu
ex_engine/xllm_kernels/cuda/norm.cu (600 lines) — rms_norm, fused_add_rms_norm
ex_engine/xllm_kernels/cuda/rope.cu (258 lines) — rotary_embedding
ex_engine/xllm_kernels/cuda/block_copy.cu (209 lines) — copy_blocks, swap_blocks
ex_engine/xllm_kernels/cuda/reshape_paged_cache.cu (101 lines) — KV cache ops
ex_engine/xllm_kernels/cuda/headers/ (5 headers for compilation)
ILU bridge kernel sources (from xllm, verified SAME as upstream):
ex_engine/xllm_kernels/ilu/ (10 files, 925 lines total)
— activation.cpp, attention.cpp, fused_moe.cpp, group_gemm.cpp,
matmul.cpp, norm.cpp, rope.cpp, ilu_ops_api.h, ixformer.h, utils.h
FLA Triton GDN kernels (for GatedDeltaNet without SM90+ FlashQLA):
ex_engine/fla_kernels/gated_delta_rule/ (7 files, 2370 lines)
— chunk_fwd.py (428), chunk.py (487), wy_fast.py (409),
fused_recurrent.py (392), naive.py (161), gate.py (380)
CCCL sync (12 tuning + 14 dispatch headers updated from NVIDIA/cccl):
cccl_upstream/cub/cub/device/dispatch/tuning/ — 12 changed files synced
cccl_upstream/cub/cub/device/dispatch/ — 14 changed dispatch files synced
Compilation targets for real machine (ivcore10):
1. CUDA kernels: --cuda-gpu-arch=ivcore10 via corex clang/16
2. ILU bridges: torch.utils.cpp_extension linking ixformer .so
3. FLA kernels: Triton JIT (if Triton works on BI-V100)
102 lines
3.9 KiB
Plaintext
102 lines
3.9 KiB
Plaintext
/* Copyright 2025-2026 The xLLM Authors. All Rights Reserved.
|
|
|
|
Licensed under the Apache License, Version 2.0 (the "License");
|
|
you may not use this file except in compliance with the License.
|
|
You may obtain a copy of the License at
|
|
|
|
https://github.com/jd-opensource/xllm/blob/main/LICENSE
|
|
|
|
Unless required by applicable law or agreed to in writing, software
|
|
distributed under the License is distributed on an "AS IS" BASIS,
|
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
See the License for the specific language governing permissions and
|
|
limitations under the License.
|
|
==============================================================================*/
|
|
|
|
#include <c10/cuda/CUDAStream.h>
|
|
|
|
#include "cuda_ops_api.h"
|
|
#include "device_utils.cuh"
|
|
|
|
namespace xllm::kernel::cuda {
|
|
|
|
template <typename T>
|
|
__global__ void XLLM_KERNEL_ATTR(1024) reshape_paged_cache_kernel(
|
|
const int* __restrict__ slot_ids, // [n_tokens]
|
|
const T* __restrict__ keys, // [n_tokens, n_heads, head_dim]
|
|
const T* __restrict__ values, // [n_tokens, n_heads, head_dim]
|
|
T* __restrict__ key_cache,
|
|
T* __restrict__ value_cache,
|
|
int64_t k_stride,
|
|
int64_t v_stride,
|
|
int64_t n_kv_heads,
|
|
int64_t head_dim,
|
|
int64_t block_size) {
|
|
// block/token index
|
|
const int64_t bid = blockIdx.x;
|
|
// which slot to write to
|
|
const int64_t slot_id = slot_ids[bid];
|
|
if (slot_id < 0) {
|
|
return;
|
|
}
|
|
// block index
|
|
const int64_t block_idx = slot_id / block_size;
|
|
// offset within block
|
|
const int64_t block_offset = slot_id % block_size;
|
|
// base index for the block in cache
|
|
const int64_t block_base_idx = block_idx * block_size * n_kv_heads * head_dim;
|
|
// copy value one by one for the token
|
|
for (int64_t i = threadIdx.x; i < n_kv_heads * head_dim; i += blockDim.x) {
|
|
const int64_t k_src_idx = bid * k_stride + i;
|
|
const int64_t v_src_idx = bid * v_stride + i;
|
|
// cache: [n_blocks, block_size, n_heads, head_dim]
|
|
const int64_t head_base_idx =
|
|
block_base_idx + block_offset * n_kv_heads * head_dim;
|
|
// which head to write to
|
|
const int head_idx = i / head_dim;
|
|
// which dim within head to write to
|
|
const int head_offset = i % head_dim;
|
|
const int64_t dst_idx = head_base_idx + head_idx * head_dim + head_offset;
|
|
key_cache[dst_idx] = keys[k_src_idx];
|
|
value_cache[dst_idx] = values[v_src_idx];
|
|
}
|
|
}
|
|
|
|
void reshape_paged_cache(
|
|
torch::Tensor slot_ids, // [n_tokens]
|
|
torch::Tensor keys, // [n_tokens, n_kv_heads, head_dim]
|
|
torch::Tensor values, // [n_tokens, n_kv_heads, head_dim]
|
|
torch::Tensor key_cache, // [n_blocks, block_size, n_heads, head_dim]
|
|
torch::Tensor value_cache) {
|
|
// keys and values should be continuous at n_kv_heads and head_dim dims
|
|
CHECK(keys.stride(-1) == 1 && keys.stride(-2) == keys.size(-1));
|
|
CHECK(values.stride(-1) == 1 && values.stride(-2) == values.size(-1));
|
|
const int64_t n_tokens = keys.size(-3);
|
|
const int64_t n_kv_heads = keys.size(-2);
|
|
const int64_t head_dim = keys.size(-1);
|
|
const int64_t block_size = key_cache.size(-3);
|
|
// it is possible that keys and values have different strides
|
|
const int64_t k_stride = keys.stride(-3);
|
|
const int64_t v_stride = values.stride(-3);
|
|
const int64_t n = n_kv_heads * head_dim;
|
|
dim3 grid(n_tokens);
|
|
dim3 block(std::min<int>(n, 1024));
|
|
DISPATCH_FLOATING_TYPES(
|
|
keys.scalar_type(), "reshape_paged_cache_kernel", [&] {
|
|
reshape_paged_cache_kernel<scalar_t>
|
|
<<<grid, block, 0, c10::cuda::getCurrentCUDAStream()>>>(
|
|
slot_ids.data_ptr<int>(),
|
|
keys.data_ptr<scalar_t>(),
|
|
values.data_ptr<scalar_t>(),
|
|
key_cache.data_ptr<scalar_t>(),
|
|
value_cache.data_ptr<scalar_t>(),
|
|
k_stride,
|
|
v_stride,
|
|
n_kv_heads,
|
|
head_dim,
|
|
block_size);
|
|
});
|
|
}
|
|
|
|
} // namespace xllm::kernel::cuda
|