Skip to content

ggml-cuda : optimize mmvq nwarps for Pascal DP4A #25479

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
animeshsri14 wants to merge 1 commit into
base: master
Choose a base branch
from
+14 −3
Draft

ggml-cuda : optimize mmvq nwarps for Pascal DP4A #25479

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
17 changes: 14 additions & 3 deletions ggml/src/ggml-cuda/mmvq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) {

enum mmvq_parameter_table_id {
MMVQ_PARAMETERS_GENERIC = 0,
MMVQ_PARAMETERS_PASCAL_DP4A,
MMVQ_PARAMETERS_TURING,
MMVQ_PARAMETERS_GCN,
MMVQ_PARAMETERS_RDNA2,
Expand All @@ -81,6 +82,8 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
return MMVQ_PARAMETERS_GCN;
#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING && __CUDA_ARCH__ < GGML_CUDA_CC_AMPERE
return MMVQ_PARAMETERS_TURING;
#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A && __CUDA_ARCH__ < GGML_CUDA_CC_VOLTA
return MMVQ_PARAMETERS_PASCAL_DP4A;
#else
return MMVQ_PARAMETERS_GENERIC;
#endif
Expand All @@ -99,9 +102,13 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) {
return MMVQ_PARAMETERS_GCN;
}
if (GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING && ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_AMPERE) {
const int arch = ggml_cuda_highest_compiled_arch(cc);
if (GGML_CUDA_CC_IS_NVIDIA(cc) && arch >= GGML_CUDA_CC_TURING && arch < GGML_CUDA_CC_AMPERE) {
return MMVQ_PARAMETERS_TURING;
}
if (GGML_CUDA_CC_IS_NVIDIA(cc) && arch >= GGML_CUDA_CC_DP4A && arch < GGML_CUDA_CC_VOLTA) {
return MMVQ_PARAMETERS_PASCAL_DP4A;
}
return MMVQ_PARAMETERS_GENERIC;
}

Expand Down Expand Up @@ -350,7 +357,11 @@ static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() {
}

static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id) {
if (table_id == MMVQ_PARAMETERS_GENERIC) {
if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_PASCAL_DP4A) {
// Pascal cc 6.1/6.2 decode is bandwidth-bound on small SMs: 2 warps beat 4 for single-token.
if (table_id == MMVQ_PARAMETERS_PASCAL_DP4A && ncols_dst == 1) {
return 2;
}
switch (ncols_dst) {
case 1:
case 2:
Expand Down Expand Up @@ -456,7 +467,7 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
}

static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int table_id, bool small_k = false, int nwarps = 1) {
if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING) {
if (table_id == MMVQ_PARAMETERS_GENERIC || table_id == MMVQ_PARAMETERS_PASCAL_DP4A || table_id == MMVQ_PARAMETERS_GCN || table_id == MMVQ_PARAMETERS_TURING) {
switch (ncols_dst) {
case 1:
return small_k ? nwarps : 1;
Expand Down