From 31d3ee99bb0361c461d2a2cefa521f629baf9563 Mon Sep 17 00:00:00 2001 From: claude Date: Fri, 14 Aug 2026 11:31:54 +0000 Subject: [PATCH] =?UTF-8?q?fix:=20MoE=20kernel=20include=20paths=20?= =?UTF-8?q?=E2=80=94=20device=5Futils.cuh=20+=20arch=5Fcondition.h?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fixed xllm internal paths to our headers/ directory: kernels/cuda/device_utils.cuh → device_utils.cuh core/kernels/cuda/device_utils.cuh → device_utils.cuh core/kernels/cuda/arch_condition.h → arch_condition.h (copied) --- .../cuda/headers/arch_condition.h | 108 ++++++++++++++++++ ex_engine/xllm_kernels/cuda/moe/moe_topk.cuh | 4 +- .../cuda/moe/moe_topk_sigmoid_kernels.cuh | 2 +- .../cuda/moe/moe_topk_softmax_kernels.cuh | 2 +- 4 files changed, 112 insertions(+), 4 deletions(-) create mode 100644 ex_engine/xllm_kernels/cuda/headers/arch_condition.h diff --git a/ex_engine/xllm_kernels/cuda/headers/arch_condition.h b/ex_engine/xllm_kernels/cuda/headers/arch_condition.h new file mode 100644 index 00000000..a424f18e --- /dev/null +++ b/ex_engine/xllm_kernels/cuda/headers/arch_condition.h @@ -0,0 +1,108 @@ +/* + * Copyright (c) 2022-2025, NVIDIA CORPORATION. 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 + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * 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. + */ + +// refers to +// https://github.com/NVIDIA/TensorRT-LLM/blob/main/cpp/include/tensorrt_llm/kernels/archCondition.h + +#pragma once + +namespace xllm::kernel::cuda { +namespace detail { + +#ifdef __CUDA_ARCH__ + +// __CUDA_ARCH_SPECIFIC__ is only available starting from CUDA 12.9 +#if (__CUDACC_VER_MAJOR__ > 12 || \ + (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ >= 9)) +#define HAS_CUDA_SPECIFIC_MACRO 1 + +#if __CUDA_ARCH__ >= 900 +#if !defined(__CUDA_ARCH_SPECIFIC__) && !defined(__CUDA_ARCH_FAMILY_SPECIFIC__) +#error \ + "Compiling for SM90 or newer architectures must use Arch specific or Arch Family specific target" +#endif +#endif + +#else +#define HAS_CUDA_SPECIFIC_MACRO 0 +#endif + +// For CUDA < 12.9, we assume that sm90 or newer architectures are always built +// with arch specific. +#if defined(__CUDA_ARCH_SPECIFIC__) || \ + (!HAS_CUDA_SPECIFIC_MACRO && __CUDA_ARCH__ >= 900) +static constexpr bool isArchSpecific = true; +#else +static constexpr bool isArchSpecific = false; +#endif + +struct arch_info { + static constexpr bool mIsDevice = true; + static constexpr bool mArchSpecific = isArchSpecific; + static constexpr int mMajor = __CUDA_ARCH__ / 100; + static constexpr int mMinor = __CUDA_ARCH__ / 10 % 10; + static constexpr int mArch = __CUDA_ARCH__ / 10; +}; + +#else + +struct arch_info { + static constexpr bool mIsDevice = false; + static constexpr bool mArchSpecific = false; + static constexpr int mMajor = 0; + static constexpr int mMinor = 0; + static constexpr int mArch = 0; +}; + +#endif + +} // namespace detail + +namespace arch { + +struct is_device : std::bool_constant {}; + +struct is_arch_specific : std::bool_constant { +}; + +template +struct is_match + : std::bool_constant { +}; + +template +struct is_major : std::bool_constant {}; + +template +struct is_compatible : std::bool_constant::value && + detail::arch_info::mArch >= Arch> {}; + +inline constexpr bool is_device_v = is_device::value; + +inline constexpr bool is_arch_specific_v = is_arch_specific::value; + +template +inline constexpr bool is_match_v = is_match::value; + +template +inline constexpr bool is_major_v = is_major::value; + +template +inline constexpr bool is_compatible_v = is_compatible::value; + +} // namespace arch +} // namespace xllm::kernel::cuda diff --git a/ex_engine/xllm_kernels/cuda/moe/moe_topk.cuh b/ex_engine/xllm_kernels/cuda/moe/moe_topk.cuh index 90e9177a..6c85d9bc 100644 --- a/ex_engine/xllm_kernels/cuda/moe/moe_topk.cuh +++ b/ex_engine/xllm_kernels/cuda/moe/moe_topk.cuh @@ -35,14 +35,14 @@ #include #endif -#include "core/kernels/cuda/arch_condition.h" +#include "arch_condition.h" #if defined(USE_DCU) #include #include #endif -#include "core/kernels/cuda/device_utils.cuh" +#include "device_utils.cuh" namespace xllm::kernel::cuda { namespace reduce_topk { diff --git a/ex_engine/xllm_kernels/cuda/moe/moe_topk_sigmoid_kernels.cuh b/ex_engine/xllm_kernels/cuda/moe/moe_topk_sigmoid_kernels.cuh index 15b4dc90..b64299cb 100644 --- a/ex_engine/xllm_kernels/cuda/moe/moe_topk_sigmoid_kernels.cuh +++ b/ex_engine/xllm_kernels/cuda/moe/moe_topk_sigmoid_kernels.cuh @@ -26,7 +26,7 @@ limitations under the License. #if !defined(USE_DCU) && !defined(USE_MACA) #endif -#include "kernels/cuda/device_utils.cuh" +#include "device_utils.cuh" namespace { diff --git a/ex_engine/xllm_kernels/cuda/moe/moe_topk_softmax_kernels.cuh b/ex_engine/xllm_kernels/cuda/moe/moe_topk_softmax_kernels.cuh index 43119a37..a552dc8e 100644 --- a/ex_engine/xllm_kernels/cuda/moe/moe_topk_softmax_kernels.cuh +++ b/ex_engine/xllm_kernels/cuda/moe/moe_topk_softmax_kernels.cuh @@ -26,7 +26,7 @@ limitations under the License. #if !defined(USE_DCU) && !defined(USE_MACA) #endif -#include "kernels/cuda/device_utils.cuh" +#include "device_utils.cuh" using cub_kvp = cub::KeyValuePair;