Skip to content

cuda/hip: optional wave64 for the quantized mat-vec kernels on RDNA #89

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
mgehre-amd wants to merge 1 commit into
base: matthias.q4k-mmvq-gfx1151
Choose a base branch
from
Draft

cuda/hip: optional wave64 for the quantized mat-vec kernels on RDNA #89

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 8 additions & 1 deletion ggml/src/ggml-cuda/common.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -381,8 +381,15 @@ struct block_q8_1_x4 {
};
static_assert(sizeof(block_q8_1_x4) == 4*sizeof(block_q8_1), "block_q8_1_x4 must alias 4 q8_1 blocks");

// Set by ggml-hip/CMakeLists.txt on the translation units built with -mwavefrontsize64.
#ifndef GGML_CUDA_FORCE_WAVE64
#define GGML_CUDA_FORCE_WAVE64 0
#endif

static constexpr __device__ int ggml_cuda_get_physical_warp_size() {
#if defined(GGML_USE_HIP) && (defined(__GFX9__) || defined(__GFX8__))
#if GGML_CUDA_FORCE_WAVE64
return 64;
#elif defined(GGML_USE_HIP) && (defined(__GFX9__) || defined(__GFX8__))
return 64;
#else
return 32;
Expand Down
5 changes: 4 additions & 1 deletion ggml/src/ggml-cuda/mmvq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -898,7 +898,10 @@ static void mul_mat_vec_q_switch_ncols_dst(

const int device = ggml_cuda_get_device();
const int cc = ggml_cuda_info().devices[device].cc;
const int warp_size = ggml_cuda_info().devices[device].warp_size;
// This TU is compiled -mwavefrontsize64 when GGML_HIP_MMVQ_WAVE64 is set, in which case its
// kernels have 64 lanes per wave regardless of what the device properties report (32 on
// RDNA), and the launch geometry has to match.
const int warp_size = GGML_CUDA_FORCE_WAVE64 ? 64 : ggml_cuda_info().devices[device].warp_size;
const mmvq_parameter_table_id table_id = get_device_table_id(cc);

const bool has_ids = ids != nullptr;
Expand Down
9 changes: 9 additions & 0 deletions ggml/src/ggml-hip/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,6 +61,15 @@ file(GLOB GGML_HEADERS_ROCM "../ggml-cuda/*.cuh")
list(APPEND GGML_HEADERS_ROCM "../../include/ggml-cuda.h")

file(GLOB GGML_SOURCES_ROCM "../ggml-cuda/*.cu")

# Opt-in: build the quantized mat-vec kernels for wave64. RDNA selects the wave size per kernel,
# but -mwavefrontsize64 is per translation unit, so this flips every kernel in mmvq.cu, not just
# the q4_K one. Not applied target-wide: fattn-mma-f16.cuh static_asserts on wave32 tiling.
if (GGML_HIP_MMVQ_WAVE64)
set_source_files_properties("../ggml-cuda/mmvq.cu" PROPERTIES
COMPILE_OPTIONS "-mwavefrontsize64;-DGGML_CUDA_FORCE_WAVE64=1")
endif()

file(GLOB SRCS "../ggml-cuda/template-instances/fattn-tile*.cu")
list(APPEND GGML_SOURCES_ROCM ${SRCS})
file(GLOB SRCS "../ggml-cuda/template-instances/fattn-mma*.cu")
Expand Down
Loading