diff --git a/CMakeLists.txt b/CMakeLists.txt index c03360a5d4e..6f60759550b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -385,76 +385,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA") endif() endif() - # - # Machete kernels - - # The machete kernels only work on hopper and require CUDA 12.0 or later. - # Only build Machete kernels if we are building for something compatible with sm90a - cuda_archs_loose_intersection(MACHETE_ARCHS "9.0a" "${CUDA_ARCHS}") - if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND MACHETE_ARCHS) - # - # For the Machete kernels we automatically generate sources for various - # preselected input type pairs and schedules. - # Generate sources: - set(MACHETE_GEN_SCRIPT - ${CMAKE_CURRENT_SOURCE_DIR}/csrc/quantization/machete/generate.py) - file(MD5 ${MACHETE_GEN_SCRIPT} MACHETE_GEN_SCRIPT_HASH) - - message(STATUS "Machete generation script hash: ${MACHETE_GEN_SCRIPT_HASH}") - message(STATUS "Last run machete generate script hash: $CACHE{MACHETE_GEN_SCRIPT_HASH}") - - if (NOT DEFINED CACHE{MACHETE_GEN_SCRIPT_HASH} - OR NOT $CACHE{MACHETE_GEN_SCRIPT_HASH} STREQUAL ${MACHETE_GEN_SCRIPT_HASH}) - execute_process( - COMMAND ${CMAKE_COMMAND} -E env - PYTHONPATH=${CMAKE_CURRENT_SOURCE_DIR}/csrc/cutlass_extensions/:${CUTLASS_DIR}/python/:${VLLM_PYTHON_PATH}:$ENV{PYTHONPATH} - ${Python_EXECUTABLE} ${MACHETE_GEN_SCRIPT} - RESULT_VARIABLE machete_generation_result - OUTPUT_VARIABLE machete_generation_output - OUTPUT_FILE ${CMAKE_CURRENT_BINARY_DIR}/machete_generation.log - ERROR_FILE ${CMAKE_CURRENT_BINARY_DIR}/machete_generation.log - ) - - if (NOT machete_generation_result EQUAL 0) - message(FATAL_ERROR "Machete generation failed." - " Result: \"${machete_generation_result}\"" - "\nCheck the log for details: " - "${CMAKE_CURRENT_BINARY_DIR}/machete_generation.log") - else() - set(MACHETE_GEN_SCRIPT_HASH ${MACHETE_GEN_SCRIPT_HASH} - CACHE STRING "Last run machete generate script hash" FORCE) - message(STATUS "Machete generation completed successfully.") - endif() - else() - message(STATUS "Machete generation script has not changed, skipping generation.") - endif() - - # Add machete generated sources - file(GLOB MACHETE_GEN_SOURCES "csrc/quantization/machete/generated/*.cu") - list(APPEND VLLM_EXT_SRC ${MACHETE_GEN_SOURCES}) - - # forward compatible - set_gencode_flags_for_srcs( - SRCS "${MACHETE_GEN_SOURCES}" - CUDA_ARCHS "${MACHETE_ARCHS}") - - list(APPEND VLLM_EXT_SRC - csrc/quantization/machete/machete_pytorch.cu) - - message(STATUS "Building Machete kernels for archs: ${MACHETE_ARCHS}") - else() - if (NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 - AND MACHETE_ARCHS) - message(STATUS "Not building Machete kernels as CUDA Compiler version is " - "not >= 12.0, we recommend upgrading to CUDA 12.0 or " - "later if you intend on running w4a16 quantized models on " - "Hopper.") - else() - message(STATUS "Not building Machete kernels as no compatible archs " - "found in CUDA target architectures") - endif() - endif() - # if CUDA endif @@ -533,6 +463,75 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP") "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu" "csrc/libtorch_stable/minimax_reduce_rms_kernel.cu") + # + # Machete kernels + # + # The machete kernels only work on hopper and require CUDA 12.0 or later. + # Only build Machete kernels if we are building for something compatible with sm90a + cuda_archs_loose_intersection(MACHETE_ARCHS "9.0a" "${CUDA_ARCHS}") + if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND MACHETE_ARCHS) + # + # For the Machete kernels we automatically generate sources for various + # preselected input type pairs and schedules. + # Generate sources: + set(MACHETE_GEN_SCRIPT + ${CMAKE_CURRENT_SOURCE_DIR}/csrc/libtorch_stable/quantization/machete/generate.py) + file(MD5 ${MACHETE_GEN_SCRIPT} MACHETE_GEN_SCRIPT_HASH) + + message(STATUS "Machete generation script hash: ${MACHETE_GEN_SCRIPT_HASH}") + message(STATUS "Last run machete generate script hash: $CACHE{MACHETE_GEN_SCRIPT_HASH}") + + if (NOT DEFINED CACHE{MACHETE_GEN_SCRIPT_HASH} + OR NOT $CACHE{MACHETE_GEN_SCRIPT_HASH} STREQUAL ${MACHETE_GEN_SCRIPT_HASH}) + execute_process( + COMMAND ${CMAKE_COMMAND} -E env + PYTHONPATH=${CMAKE_CURRENT_SOURCE_DIR}/csrc/cutlass_extensions/:${CUTLASS_DIR}/python/:${VLLM_PYTHON_PATH}:$ENV{PYTHONPATH} + ${Python_EXECUTABLE} ${MACHETE_GEN_SCRIPT} + RESULT_VARIABLE machete_generation_result + OUTPUT_VARIABLE machete_generation_output + OUTPUT_FILE ${CMAKE_CURRENT_BINARY_DIR}/machete_generation.log + ERROR_FILE ${CMAKE_CURRENT_BINARY_DIR}/machete_generation.log + ) + + if (NOT machete_generation_result EQUAL 0) + message(FATAL_ERROR "Machete generation failed." + " Result: \"${machete_generation_result}\"" + "\nCheck the log for details: " + "${CMAKE_CURRENT_BINARY_DIR}/machete_generation.log") + else() + set(MACHETE_GEN_SCRIPT_HASH ${MACHETE_GEN_SCRIPT_HASH} + CACHE STRING "Last run machete generate script hash" FORCE) + message(STATUS "Machete generation completed successfully.") + endif() + else() + message(STATUS "Machete generation script has not changed, skipping generation.") + endif() + + # Add machete generated sources + file(GLOB MACHETE_GEN_SOURCES "csrc/libtorch_stable/quantization/machete/generated/*.cu") + list(APPEND VLLM_STABLE_EXT_SRC ${MACHETE_GEN_SOURCES}) + + # forward compatible + set_gencode_flags_for_srcs( + SRCS "${MACHETE_GEN_SOURCES}" + CUDA_ARCHS "${MACHETE_ARCHS}") + + list(APPEND VLLM_STABLE_EXT_SRC + csrc/libtorch_stable/quantization/machete/machete_pytorch.cu) + message(STATUS "Building Machete kernels for archs: ${MACHETE_ARCHS}") + else() + if (NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 + AND MACHETE_ARCHS) + message(STATUS "Not building Machete kernels as CUDA Compiler version is " + "not >= 12.0, we recommend upgrading to CUDA 12.0 or " + "later if you intend on running w4a16 quantized models on " + "Hopper.") + else() + message(STATUS "Not building Machete kernels as no compatible archs " + "found in CUDA target architectures") + endif() + endif() + set_gencode_flags_for_srcs( SRCS "${VLLM_STABLE_EXT_SRC}" CUDA_ARCHS "${CUDA_ARCHS}") diff --git a/csrc/cutlass_extensions/vllm_cutlass_library_extension.py b/csrc/cutlass_extensions/vllm_cutlass_library_extension.py index 34fb64c413d..d692502f3ff 100644 --- a/csrc/cutlass_extensions/vllm_cutlass_library_extension.py +++ b/csrc/cutlass_extensions/vllm_cutlass_library_extension.py @@ -57,13 +57,13 @@ VLLMDataTypeVLLMScalarTypeTag: dict[VLLMDataType | DataType, str] = { } VLLMDataTypeTorchDataTypeTag: dict[VLLMDataType | DataType, str] = { - DataType.u8: "at::ScalarType::Byte", - DataType.s8: "at::ScalarType::Char", - DataType.e4m3: "at::ScalarType::Float8_e4m3fn", - DataType.s32: "at::ScalarType::Int", - DataType.f16: "at::ScalarType::Half", - DataType.bf16: "at::ScalarType::BFloat16", - DataType.f32: "at::ScalarType::Float", + DataType.u8: "torch::headeronly::ScalarType::Byte", + DataType.s8: "torch::headeronly::ScalarType::Char", + DataType.e4m3: "torch::headeronly::ScalarType::Float8_e4m3fn", + DataType.s32: "torch::headeronly::ScalarType::Int", + DataType.f16: "torch::headeronly::ScalarType::Half", + DataType.bf16: "torch::headeronly::ScalarType::BFloat16", + DataType.f32: "torch::headeronly::ScalarType::Float", } VLLMKernelScheduleTag: dict[MixedInputKernelScheduleType | KernelScheduleType, str] = { diff --git a/csrc/quantization/machete/Readme.md b/csrc/libtorch_stable/quantization/machete/Readme.md similarity index 100% rename from csrc/quantization/machete/Readme.md rename to csrc/libtorch_stable/quantization/machete/Readme.md diff --git a/csrc/quantization/machete/generate.py b/csrc/libtorch_stable/quantization/machete/generate.py similarity index 95% rename from csrc/quantization/machete/generate.py rename to csrc/libtorch_stable/quantization/machete/generate.py index e12601e9e97..11a5bbdd13c 100644 --- a/csrc/quantization/machete/generate.py +++ b/csrc/libtorch_stable/quantization/machete/generate.py @@ -39,10 +39,10 @@ namespace machete { {% for impl_config in impl_configs %} {% set type_sig = gen_type_sig(impl_config.types) -%} {% for s in impl_config.schedules %} -extern torch::Tensor impl_{{type_sig}}_sch_{{gen_sch_sig(s)}}(MMArgs); +extern torch::stable::Tensor impl_{{type_sig}}_sch_{{gen_sch_sig(s)}}(MMArgs); {%- endfor %} -torch::Tensor mm_dispatch_{{type_sig}}(MMArgs args) { +torch::stable::Tensor mm_dispatch_{{type_sig}}(MMArgs args) { [[maybe_unused]] auto M = args.A.size(0); [[maybe_unused]] auto N = args.B.size(1); [[maybe_unused]] auto K = args.A.size(1); @@ -59,14 +59,14 @@ torch::Tensor mm_dispatch_{{type_sig}}(MMArgs args) { if (*args.maybe_schedule == "{{ gen_sch_sig(s) }}") return impl_{{type_sig}}_sch_{{ gen_sch_sig(s) }}(args); {%- endfor %} - TORCH_CHECK_NOT_IMPLEMENTED(false, "machete_gemm(..) is not implemented for " + STD_TORCH_CHECK_NOT_IMPLEMENTED(false, "machete_gemm(..) is not implemented for " "schedule = ", *args.maybe_schedule); } {%- endfor %} -static inline std::optional maybe_scalartype( - std::optional const& t) { +static inline std::optional maybe_scalartype( + std::optional const& t) { if (!t) { return std::nullopt; } else { @@ -74,7 +74,7 @@ static inline std::optional maybe_scalartype( }; } -torch::Tensor mm_dispatch(MMArgs args) { +torch::stable::Tensor mm_dispatch(MMArgs args) { auto out_type = args.maybe_out_type.value_or(args.A.scalar_type()); auto a_type = args.A.scalar_type(); auto maybe_g_scales_type = maybe_scalartype(args.maybe_group_scales); @@ -105,19 +105,19 @@ torch::Tensor mm_dispatch(MMArgs args) { } {%- endfor %} - TORCH_CHECK_NOT_IMPLEMENTED( + STD_TORCH_CHECK_NOT_IMPLEMENTED( false, "machete_mm(..) is not implemented for " - "a_type=", args.A.scalar_type(), + "a_type=", torch::headeronly::toString(args.A.scalar_type()), ", b_type=", args.b_type.str(), - ", out_type=", out_type, + ", out_type=", torch::headeronly::toString(out_type), ", with_group_scale_type=", maybe_g_scales_type - ? toString(*maybe_g_scales_type) : "None", + ? torch::headeronly::toString(*maybe_g_scales_type) : "None", ", with_group_zeropoint_type=", maybe_g_zeros_type - ? toString(*maybe_g_zeros_type) : "None", + ? torch::headeronly::toString(*maybe_g_zeros_type) : "None", ", with_channel_scale_type=", maybe_ch_scales_type - ? toString(*maybe_ch_scales_type) : "None", + ? torch::headeronly::toString(*maybe_ch_scales_type) : "None", ", with_token_scale_type=", maybe_tok_scales_type - ? toString(*maybe_tok_scales_type) : "None", + ? torch::headeronly::toString(*maybe_tok_scales_type) : "None", "; implemented types are: \\n", {%- for impl_config in impl_configs %} {% set t = impl_config.types -%} @@ -197,7 +197,7 @@ using Kernel_{{type_sig}} = MacheteKernelTemplate< {% for sch in schs %} {% set sch_sig = gen_sch_sig(sch) -%} -torch::Tensor +torch::stable::Tensor impl_{{type_sig}}_sch_{{sch_sig}}(MMArgs args) { return run_impl>(args); } @@ -212,7 +212,7 @@ PREPACK_TEMPLATE = """ namespace machete { -torch::Tensor prepack_B_dispatch(PrepackBArgs args) { +torch::stable::Tensor prepack_B_dispatch(PrepackBArgs args) { auto convert_type = args.maybe_group_scales_type.value_or(args.a_type); {%- for t in types %} {% set b_type = unsigned_type_with_bitwidth(t.b_num_bits) %} @@ -231,12 +231,12 @@ torch::Tensor prepack_B_dispatch(PrepackBArgs args) { } {%- endfor %} - TORCH_CHECK_NOT_IMPLEMENTED(false, + STD_TORCH_CHECK_NOT_IMPLEMENTED(false, "prepack_B_dispatch(..) is not implemented for " - "atype = ", args.a_type, + "atype = ", torch::headeronly::toString(args.a_type), ", b_type = ", args.b_type.str(), ", with_group_scales_type= ", args.maybe_group_scales_type ? - toString(*args.maybe_group_scales_type) : "None"); + torch::headeronly::toString(*args.maybe_group_scales_type) : "None"); } }; // namespace machete diff --git a/csrc/quantization/machete/machete_collective_builder.cuh b/csrc/libtorch_stable/quantization/machete/machete_collective_builder.cuh similarity index 100% rename from csrc/quantization/machete/machete_collective_builder.cuh rename to csrc/libtorch_stable/quantization/machete/machete_collective_builder.cuh diff --git a/csrc/quantization/machete/machete_interleaving_utils.cuh b/csrc/libtorch_stable/quantization/machete/machete_interleaving_utils.cuh similarity index 100% rename from csrc/quantization/machete/machete_interleaving_utils.cuh rename to csrc/libtorch_stable/quantization/machete/machete_interleaving_utils.cuh diff --git a/csrc/quantization/machete/machete_mainloop.cuh b/csrc/libtorch_stable/quantization/machete/machete_mainloop.cuh similarity index 100% rename from csrc/quantization/machete/machete_mainloop.cuh rename to csrc/libtorch_stable/quantization/machete/machete_mainloop.cuh diff --git a/csrc/quantization/machete/machete_mm_kernel.cuh b/csrc/libtorch_stable/quantization/machete/machete_mm_kernel.cuh similarity index 87% rename from csrc/quantization/machete/machete_mm_kernel.cuh rename to csrc/libtorch_stable/quantization/machete/machete_mm_kernel.cuh index cc50e68b058..db3321a39db 100644 --- a/csrc/quantization/machete/machete_mm_kernel.cuh +++ b/csrc/libtorch_stable/quantization/machete/machete_mm_kernel.cuh @@ -1,8 +1,6 @@ #pragma once -#include -#include -#include +#include // clang-format off // The cutlass include order matters (annoyingly) @@ -175,19 +173,23 @@ struct MacheteKernelTemplate { static Arguments create_arguments( cudaStream_t stream, - torch::Tensor const& A, // MxK matrix - torch::Tensor const& B, // KxN prepacked matrix - torch::Tensor& D, // MxN matrix - std::optional const& maybe_g_scales, // scale_KxN matrix - std::optional const& maybe_g_zeros, // scale_KxN matrix + torch::stable::Tensor const& A, // MxK matrix + torch::stable::Tensor const& B, // KxN prepacked matrix + torch::stable::Tensor& D, // MxN matrix + std::optional const& + maybe_g_scales, // scale_KxN matrix + std::optional const& + maybe_g_zeros, // scale_KxN matrix std::optional maybe_group_size, - std::optional const& maybe_ch_scales, // len N vector - std::optional const& maybe_tok_scales) // len M vector + std::optional const& + maybe_ch_scales, // len N vector + std::optional const& + maybe_tok_scales) // len M vector { static_assert(!with_group_zeropoints || with_group_scales); int M = A.size(0), N = B.size(1), K = A.size(1); - TORCH_CHECK(D.size(0) == M && D.size(1) == N); + STD_TORCH_CHECK(D.size(0) == M && D.size(1) == N); auto layout_A = make_cute_layout(A, "A"); auto layout_D = make_cute_layout(D, "D"); @@ -216,29 +218,29 @@ struct MacheteKernelTemplate { maybe_group_size == -1 ? K : maybe_group_size.value_or(K); int const scale_k = (K + group_size - 1) / group_size; - TORCH_CHECK(size<0>(layout_A) == M && size<1>(layout_A) == K); - TORCH_CHECK(size<0>(layout_D) == M && size<1>(layout_D) == N); + STD_TORCH_CHECK(size<0>(layout_A) == M && size<1>(layout_A) == K); + STD_TORCH_CHECK(size<0>(layout_D) == M && size<1>(layout_D) == N); if constexpr (with_group_scales) { - TORCH_CHECK(S_group_ptr && layout_S_group); - TORCH_CHECK((size<0>(*layout_S_group) == scale_k && - size<1>(*layout_S_group) == N)); + STD_TORCH_CHECK(S_group_ptr && layout_S_group); + STD_TORCH_CHECK((size<0>(*layout_S_group) == scale_k && + size<1>(*layout_S_group) == N)); } else { - TORCH_CHECK(!S_group_ptr, "Scales not supported"); + STD_TORCH_CHECK(!S_group_ptr, "Scales not supported"); } if constexpr (with_group_zeropoints) { - TORCH_CHECK(Z_group_ptr && layout_Z_group); - TORCH_CHECK((size<0>(*layout_Z_group) == scale_k && - size<1>(*layout_Z_group) == N)); - TORCH_CHECK(layout_S_group && *layout_Z_group == *layout_S_group, - "Scales and zeros must have the same layout"); + STD_TORCH_CHECK(Z_group_ptr && layout_Z_group); + STD_TORCH_CHECK((size<0>(*layout_Z_group) == scale_k && + size<1>(*layout_Z_group) == N)); + STD_TORCH_CHECK(layout_S_group && *layout_Z_group == *layout_S_group, + "Scales and zeros must have the same layout"); } else { - TORCH_CHECK(!Z_group_ptr, "Zeropoints not supported"); + STD_TORCH_CHECK(!Z_group_ptr, "Zeropoints not supported"); } if constexpr (with_channel_scales || with_token_scales) { - TORCH_CHECK( + STD_TORCH_CHECK( (maybe_ch_scales->numel() == N || maybe_ch_scales->numel() == 1) && (maybe_tok_scales->numel() == M || maybe_tok_scales->numel() == 1)); } @@ -298,11 +300,12 @@ struct MacheteKernelTemplate { Gemm gemm_op; cutlass::Status status = gemm_op.initialize(args, workspace, stream); - TORCH_CHECK(status == cutlass::Status::kSuccess, - "Machete kernel failed to initialize workspace"); + STD_TORCH_CHECK(status == cutlass::Status::kSuccess, + "Machete kernel failed to initialize workspace"); status = gemm_op.run(stream); - TORCH_CHECK(status == cutlass::Status::kSuccess, "Machete kernel failed"); + STD_TORCH_CHECK(status == cutlass::Status::kSuccess, + "Machete kernel failed"); } }; diff --git a/csrc/libtorch_stable/quantization/machete/machete_mm_launcher.cuh b/csrc/libtorch_stable/quantization/machete/machete_mm_launcher.cuh new file mode 100644 index 00000000000..fcf7f18aac2 --- /dev/null +++ b/csrc/libtorch_stable/quantization/machete/machete_mm_launcher.cuh @@ -0,0 +1,80 @@ +#pragma once + +#include "machete_mm_kernel.cuh" +#include "cutlass_extensions/torch_utils.hpp" +#include "core/scalar_type.hpp" +#include "libtorch_stable/torch_utils.h" + +#include +#include + +#include +#include +#include + +namespace machete { + +struct MMArgs { + torch::stable::Tensor const& A; + torch::stable::Tensor const& B; + vllm::ScalarType const& b_type; + std::optional const& maybe_out_type; + std::optional const& maybe_group_scales; + std::optional const& maybe_group_zeros; + std::optional maybe_group_size; + std::optional const& maybe_channel_scales; + std::optional const& maybe_token_scales; + std::optional maybe_schedule; +}; + +struct SupportedSchedulesArgs { + torch::headeronly::ScalarType a_type; + vllm::ScalarType b_type; + std::optional maybe_group_scales_type; + std::optional maybe_group_zeros_type; + std::optional maybe_channel_scales_type; + std::optional maybe_token_scales_type; + std::optional maybe_out_type; +}; + +torch::stable::Tensor mm_dispatch(MMArgs args); + +std::vector supported_schedules_dispatch( + SupportedSchedulesArgs args); + +template +torch::stable::Tensor run_impl(MMArgs args) { + const torch::stable::accelerator::DeviceGuard device_guard( + args.A.get_device_index()); + + auto device = args.A.device(); + auto stream = get_current_cuda_stream(device.index()); + + int M = args.A.size(0); + int N = args.B.size(1); + int K = args.A.size(1); + + // Allocate output + torch::stable::Tensor D = torch::stable::empty( + {M, N}, equivalent_scalar_type_v, + std::nullopt, device); + + auto arguments = MacheteKernel::create_arguments( + stream, // + args.A, args.B, D, args.maybe_group_scales, args.maybe_group_zeros, + args.maybe_group_size, args.maybe_channel_scales, + args.maybe_token_scales); + STD_TORCH_CHECK(MacheteKernel::can_implement(arguments), + "Machete kernel cannot be run with these arguments"); + + size_t workspace_size = MacheteKernel::get_workspace_size(arguments); + torch::stable::Tensor workspace = + torch::stable::empty(workspace_size, torch::headeronly::ScalarType::Byte, + std::nullopt, device); + + MacheteKernel::run(arguments, workspace.mutable_data_ptr(), stream); + + return D; +}; + +}; // namespace machete \ No newline at end of file diff --git a/csrc/quantization/machete/machete_prepack_kernel.cuh b/csrc/libtorch_stable/quantization/machete/machete_prepack_kernel.cuh similarity index 94% rename from csrc/quantization/machete/machete_prepack_kernel.cuh rename to csrc/libtorch_stable/quantization/machete/machete_prepack_kernel.cuh index d002355ca49..e1e054e5a00 100644 --- a/csrc/quantization/machete/machete_prepack_kernel.cuh +++ b/csrc/libtorch_stable/quantization/machete/machete_prepack_kernel.cuh @@ -3,6 +3,7 @@ #include "machete_mm_kernel.cuh" #include "cutlass_extensions/cute_utils.cuh" #include "cutlass_extensions/torch_utils.hpp" +#include namespace machete { @@ -60,8 +61,8 @@ static void prepack_B_template( auto ilvd_NKbNbKL_to_offset = PrepackedLayoutB::ilvd_NKbNbKL_to_offset(shape(B_layout)); - TORCH_CHECK(size<0>(B_layout) % size<0>(TileShapeNKL{}) == 0); - TORCH_CHECK(size<1>(B_layout) % size<1>(TileShapeNKL{}) == 0); + STD_TORCH_CHECK(size<0>(B_layout) % size<0>(TileShapeNKL{}) == 0); + STD_TORCH_CHECK(size<1>(B_layout) % size<1>(TileShapeNKL{}) == 0); auto N_tiles = size<0>(B_layout) / size<0>(TileShapeNKL{}); auto K_tiles = size<1>(B_layout) / size<1>(TileShapeNKL{}); diff --git a/csrc/quantization/machete/machete_prepack_launcher.cuh b/csrc/libtorch_stable/quantization/machete/machete_prepack_launcher.cuh similarity index 65% rename from csrc/quantization/machete/machete_prepack_launcher.cuh rename to csrc/libtorch_stable/quantization/machete/machete_prepack_launcher.cuh index 634b651a4d1..94f6f684bc0 100644 --- a/csrc/quantization/machete/machete_prepack_launcher.cuh +++ b/csrc/libtorch_stable/quantization/machete/machete_prepack_launcher.cuh @@ -3,39 +3,47 @@ #include "machete_prepack_kernel.cuh" #include "cutlass_extensions/torch_utils.hpp" #include "core/scalar_type.hpp" +#include "libtorch_stable/torch_utils.h" + +#include +#include + +#include namespace machete { struct PrepackBArgs { - torch::Tensor const& B; - at::ScalarType a_type; + torch::stable::Tensor const& B; + torch::headeronly::ScalarType a_type; vllm::ScalarType b_type; - std::optional maybe_group_scales_type; + std::optional maybe_group_scales_type; }; template -torch::Tensor prepack_impl(torch::Tensor const B) { - const at::cuda::OptionalCUDAGuard device_guard(device_of(B)); +torch::stable::Tensor prepack_impl(torch::stable::Tensor const& B) { + const torch::stable::accelerator::DeviceGuard device_guard( + B.get_device_index()); using ElementB = typename PrepackedLayoutB::ElementB; using PPBlockShape_NK = typename PrepackedLayoutB::PPBlockShape_NK; auto device = B.device(); - auto stream = at::cuda::getCurrentCUDAStream(device.index()); + auto stream = get_current_cuda_stream(device.index()); auto B_ptr = static_cast(B.const_data_ptr()); // elements per storage item for B auto eles_per_storage = - (B.dtype().itemsize() * 8) / cute::sizeof_bits_v; + (B.element_size() * 8) / cute::sizeof_bits_v; // torch B passed in is/should be (packed_K,N), the kernel expects (N,K,L) (to // match cutlass using (N,K,L) for B), so we transpose B to (N,packed_K,L) - auto Bt_packed = B.t(); + auto Bt_packed = torch::stable::transpose(B, 0, 1); - TORCH_CHECK( + STD_TORCH_CHECK( (B.size(0) * eles_per_storage) % size<1>(PPBlockShape_NK{}) == 0, "B.shape[0] (in terms of unpacked elements) must be a multiple of ", size<1>(PPBlockShape_NK{})); - TORCH_CHECK(B.size(1) % size<0>(PPBlockShape_NK{}) == 0, - "B.shape[1] must be a multiple of ", size<0>(PPBlockShape_NK{})); + STD_TORCH_CHECK(B.size(1) % size<0>(PPBlockShape_NK{}) == 0, + "B.shape[1] must be a multiple of ", + size<0>(PPBlockShape_NK{})); using StrideB = cutlass::detail::TagToStrideB_t; auto const l_Bt_packed = make_cute_layout(Bt_packed, "B"); @@ -49,7 +57,7 @@ torch::Tensor prepack_impl(torch::Tensor const B) { // new_shape = (N, packed_K, L) * (1, eles_per_storage, 1) -> (N, K, L) // new_stride = (s0, s1, s2) * (eles_per_storage, 1, eles_per_storage) // when s1 == 1 - TORCH_CHECK(stride<1>(l_Bt_packed) == 1); + STD_TORCH_CHECK(stride<1>(l_Bt_packed) == 1); // clang-format off auto const layout_Bt = make_layout( transform_with_idx(l_Bt_packed.shape(), [&](auto ele, auto idx) { @@ -61,7 +69,9 @@ torch::Tensor prepack_impl(torch::Tensor const B) { // clang-format on // Allocate output - torch::Tensor D = torch::empty_like(B, {}, at::MemoryFormat::Contiguous); + torch::stable::Tensor D = torch::stable::empty( + B.sizes(), B.scalar_type(), std::nullopt, B.device(), std::nullopt, + torch::headeronly::MemoryFormat::Contiguous); prepack_B_template( stream, B_ptr, layout_Bt, static_cast(D.mutable_data_ptr())); @@ -69,6 +79,6 @@ torch::Tensor prepack_impl(torch::Tensor const B) { return D; }; -torch::Tensor prepack_B_dispatch(PrepackBArgs args); +torch::stable::Tensor prepack_B_dispatch(PrepackBArgs args); }; // namespace machete \ No newline at end of file diff --git a/csrc/quantization/machete/machete_prepacked_layout.cuh b/csrc/libtorch_stable/quantization/machete/machete_prepacked_layout.cuh similarity index 99% rename from csrc/quantization/machete/machete_prepacked_layout.cuh rename to csrc/libtorch_stable/quantization/machete/machete_prepacked_layout.cuh index 4a7d6341e6c..c16a2ab8a33 100644 --- a/csrc/quantization/machete/machete_prepacked_layout.cuh +++ b/csrc/libtorch_stable/quantization/machete/machete_prepacked_layout.cuh @@ -1,9 +1,5 @@ #pragma once -#include -#include -#include - // clang-format off // The cutlass include order matters (annoyingly) diff --git a/csrc/libtorch_stable/quantization/machete/machete_pytorch.cu b/csrc/libtorch_stable/quantization/machete/machete_pytorch.cu new file mode 100644 index 00000000000..7736d5b3ece --- /dev/null +++ b/csrc/libtorch_stable/quantization/machete/machete_pytorch.cu @@ -0,0 +1,77 @@ +#include "machete_mm_launcher.cuh" +#include "machete_prepack_launcher.cuh" +#include "core/scalar_type.hpp" + +#include +#include +#include + +namespace machete { + +using namespace vllm; + +std::vector supported_schedules( + torch::headeronly::ScalarType a_type, int64_t b_type_id, + std::optional maybe_group_scales_type, + std::optional maybe_group_zeros_type, + std::optional maybe_channel_scales_type, + std::optional maybe_token_scales_type, + std::optional maybe_out_type) { + ScalarType const b_type = ScalarType::from_id(b_type_id); + return supported_schedules_dispatch({ + .a_type = a_type, + .b_type = b_type, + .maybe_group_scales_type = maybe_group_scales_type, + .maybe_group_zeros_type = maybe_group_zeros_type, + .maybe_channel_scales_type = maybe_channel_scales_type, + .maybe_token_scales_type = maybe_token_scales_type, + .maybe_out_type = maybe_out_type, + }); +} + +torch::stable::Tensor mm( + torch::stable::Tensor const& A, torch::stable::Tensor const& B, + int64_t b_type_id, + std::optional const& maybe_out_type, + std::optional const& maybe_group_scales, + std::optional const& maybe_group_zeros, + std::optional maybe_group_size, + std::optional const& maybe_channel_scales, + std::optional const& maybe_token_scales, + std::optional maybe_schedule) { + ScalarType const b_type = ScalarType::from_id(b_type_id); + return mm_dispatch({.A = A, + .B = B, + .b_type = b_type, + .maybe_out_type = maybe_out_type, + .maybe_group_scales = maybe_group_scales, + .maybe_group_zeros = maybe_group_zeros, + .maybe_group_size = maybe_group_size, + .maybe_channel_scales = maybe_channel_scales, + .maybe_token_scales = maybe_token_scales, + .maybe_schedule = maybe_schedule}); +} + +torch::stable::Tensor prepack_B( + torch::stable::Tensor const& B, torch::headeronly::ScalarType const& a_type, + int64_t b_type_id, + std::optional const& + maybe_group_scales_type) { + ScalarType const b_type = ScalarType::from_id(b_type_id); + return prepack_B_dispatch( + {.B = B, + .a_type = a_type, + .b_type = b_type, + .maybe_group_scales_type = maybe_group_scales_type}); +} + +STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) { + m.impl("machete_prepack_B", TORCH_BOX(&prepack_B)); + m.impl("machete_mm", TORCH_BOX(&mm)); +} + +STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, m) { + m.impl("machete_supported_schedules", TORCH_BOX(&supported_schedules)); +} + +}; // namespace machete diff --git a/csrc/libtorch_stable/torch_bindings.cpp b/csrc/libtorch_stable/torch_bindings.cpp index 204feed4a25..c805ecba1ba 100644 --- a/csrc/libtorch_stable/torch_bindings.cpp +++ b/csrc/libtorch_stable/torch_bindings.cpp @@ -34,6 +34,39 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) { // TODO: Remove this once ROCm upgrade to torch 2.11. ops.def("get_cuda_view_from_cpu_tensor(Tensor cpu_tensor) -> Tensor"); + // Machete (Dense) Optimized Mixed Precision GEMM for Hopper. + ops.def( + "machete_supported_schedules(" + " ScalarType a_type," + " int b_type," + " ScalarType? maybe_group_scales_type," + " ScalarType? maybe_group_zeros_type," + " ScalarType? maybe_channel_scales_type," + " ScalarType? maybe_token_scales_type," + " ScalarType? maybe_out_type" + ") -> str[]"); + ops.def( + "machete_mm(" + " Tensor A," + " Tensor B," + " int b_type," + " ScalarType? out_type," + " Tensor? group_scales," + " Tensor? group_zeros," + " int? group_size," + " Tensor? channel_scales," + " Tensor? token_scales," + " str? schedule" + ") -> Tensor"); + ops.def( + "machete_prepack_B(" + " Tensor B," + " ScalarType a_type," + " int b_type," + " ScalarType? group_scales_type" + ") -> Tensor"); + // conditionally compiled so impl registration is in source file + // Marlin GEMM ops.def( "marlin_gemm(Tensor a, Tensor? c_or_none, Tensor b_q_weight, " diff --git a/csrc/quantization/machete/machete_mm_launcher.cuh b/csrc/quantization/machete/machete_mm_launcher.cuh deleted file mode 100644 index cabe0af46f0..00000000000 --- a/csrc/quantization/machete/machete_mm_launcher.cuh +++ /dev/null @@ -1,75 +0,0 @@ -#pragma once - -#include -#include - -#include "machete_mm_kernel.cuh" -#include "cutlass_extensions/torch_utils.hpp" -#include "core/scalar_type.hpp" - -namespace machete { - -struct MMArgs { - torch::Tensor const& A; - torch::Tensor const& B; - vllm::ScalarType const& b_type; - std::optional const& maybe_out_type; - std::optional const& maybe_group_scales; - std::optional const& maybe_group_zeros; - std::optional maybe_group_size; - std::optional const& maybe_channel_scales; - std::optional const& maybe_token_scales; - std::optional maybe_schedule; -}; - -struct SupportedSchedulesArgs { - at::ScalarType a_type; - vllm::ScalarType b_type; - std::optional maybe_group_scales_type; - std::optional maybe_group_zeros_type; - std::optional maybe_channel_scales_type; - std::optional maybe_token_scales_type; - std::optional maybe_out_type; -}; - -torch::Tensor mm_dispatch(MMArgs args); - -std::vector supported_schedules_dispatch( - SupportedSchedulesArgs args); - -template -torch::Tensor run_impl(MMArgs args) { - const at::cuda::OptionalCUDAGuard device_guard(device_of(args.A)); - - auto device = args.A.device(); - auto stream = at::cuda::getCurrentCUDAStream(device.index()); - - int M = args.A.size(0); - int N = args.B.size(1); - int K = args.A.size(1); - - // Allocate output - torch::Tensor D = torch::empty( - {M, N}, - torch::TensorOptions() - .dtype(equivalent_scalar_type_v) - .device(device)); - - auto arguments = MacheteKernel::create_arguments( - stream, // - args.A, args.B, D, args.maybe_group_scales, args.maybe_group_zeros, - args.maybe_group_size, args.maybe_channel_scales, - args.maybe_token_scales); - TORCH_CHECK(MacheteKernel::can_implement(arguments), - "Machete kernel cannot be run with these arguments"); - - size_t workspace_size = MacheteKernel::get_workspace_size(arguments); - torch::Tensor workspace = torch::empty( - workspace_size, torch::TensorOptions().dtype(torch::kU8).device(device)); - - MacheteKernel::run(arguments, workspace.mutable_data_ptr(), stream); - - return D; -}; - -}; // namespace machete \ No newline at end of file diff --git a/csrc/quantization/machete/machete_pytorch.cu b/csrc/quantization/machete/machete_pytorch.cu deleted file mode 100644 index 05a51ee21dd..00000000000 --- a/csrc/quantization/machete/machete_pytorch.cu +++ /dev/null @@ -1,73 +0,0 @@ -#include "machete_mm_launcher.cuh" -#include "machete_prepack_launcher.cuh" -#include "core/scalar_type.hpp" - -#include "core/registration.h" - -namespace machete { - -using namespace vllm; - -std::vector supported_schedules( - at::ScalarType a_type, int64_t b_type_id, - std::optional maybe_group_scales_type, - std::optional maybe_group_zeros_type, - std::optional maybe_channel_scales_type, - std::optional maybe_token_scales_type, - std::optional maybe_out_type) { - ScalarType const b_type = ScalarType::from_id(b_type_id); - return supported_schedules_dispatch({ - .a_type = a_type, - .b_type = b_type, - .maybe_group_scales_type = maybe_group_scales_type, - .maybe_group_zeros_type = maybe_group_zeros_type, - .maybe_channel_scales_type = maybe_channel_scales_type, - .maybe_token_scales_type = maybe_token_scales_type, - .maybe_out_type = maybe_out_type, - }); -} - -torch::Tensor mm(torch::Tensor const& A, torch::Tensor const& B, - int64_t b_type_id, - std::optional const& maybe_out_type, - std::optional const& maybe_group_scales, - std::optional const& maybe_group_zeros, - std::optional maybe_group_size, - std::optional const& maybe_channel_scales, - std::optional const& maybe_token_scales, - std::optional maybe_schedule) { - ScalarType const b_type = ScalarType::from_id(b_type_id); - return mm_dispatch({.A = A, - .B = B, - .b_type = b_type, - .maybe_out_type = maybe_out_type, - .maybe_group_scales = maybe_group_scales, - .maybe_group_zeros = maybe_group_zeros, - .maybe_group_size = maybe_group_size, - .maybe_channel_scales = maybe_channel_scales, - .maybe_token_scales = maybe_token_scales, - .maybe_schedule = maybe_schedule}); -} - -torch::Tensor prepack_B( - torch::Tensor const& B, at::ScalarType const& a_type, int64_t b_type_id, - std::optional const& maybe_group_scales_type) { - ScalarType const b_type = ScalarType::from_id(b_type_id); - return prepack_B_dispatch( - {.B = B, - .a_type = a_type, - .b_type = b_type, - .maybe_group_scales_type = maybe_group_scales_type}); -} - -TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) { - m.impl("machete_prepack_B", &prepack_B); - m.impl("machete_mm", &mm); -} - -// use CatchAll since supported_schedules has no tensor arguments -TORCH_LIBRARY_IMPL(TORCH_EXTENSION_NAME, CatchAll, m) { - m.impl("machete_supported_schedules", &supported_schedules); -} - -}; // namespace machete diff --git a/csrc/torch_bindings.cpp b/csrc/torch_bindings.cpp index 941e4a61c1a..cfd185394a4 100644 --- a/csrc/torch_bindings.cpp +++ b/csrc/torch_bindings.cpp @@ -68,39 +68,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) { // custom types: // https://docs.google.com/document/d/18fBMPuOJ0fY5ZQ6YyrHUppw9FA332CpNtgB6SOIgyuA - // Machete (Dense) Optimized Mixed Precision GEMM for Hopper. - ops.def( - "machete_supported_schedules(" - " ScalarType a_type," - " int b_type," - " ScalarType? maybe_group_scales_type," - " ScalarType? maybe_group_zeros_type," - " ScalarType? maybe_channel_scales_type," - " ScalarType? maybe_token_scales_type," - " ScalarType? maybe_out_type" - ") -> str[]"); - ops.def( - "machete_mm(" - " Tensor A," - " Tensor B," - " int b_type," - " ScalarType? out_type," - " Tensor? group_scales," - " Tensor? group_zeros," - " int? group_size," - " Tensor? channel_scales," - " Tensor? token_scales," - " str? schedule" - ") -> Tensor"); - ops.def( - "machete_prepack_B(" - " Tensor B," - " ScalarType a_type," - " int b_type," - " ScalarType? group_scales_type" - ") -> Tensor"); - // conditionally compiled so impl registration is in source file - #endif }