From d55e58410e0cf95d26f959d2c086297ac9b0e01b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Sun, 9 Feb 2025 19:30:15 +0100 Subject: [PATCH 1/4] CUDA: use arch list for feature availability check --- ggml/src/ggml-common.h | 2 - ggml/src/ggml-cuda/common.cuh | 72 ++++++++++++++++++++++++++++++--- ggml/src/ggml-cuda/convert.cu | 2 +- ggml/src/ggml-cuda/ggml-cuda.cu | 3 +- ggml/src/ggml-cuda/mmq.cu | 7 ++-- ggml/src/ggml-cuda/mmq.cuh | 14 ++++--- 6 files changed, 80 insertions(+), 20 deletions(-) diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index f13fd4dea6f41..6c02b69ea239a 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -473,7 +473,6 @@ GGML_TABLE_BEGIN(uint8_t, ksigns_iq2xs, 128) 240, 113, 114, 243, 116, 245, 246, 119, 120, 249, 250, 123, 252, 125, 126, 255, GGML_TABLE_END() -//#if __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A // lowest compute capability for integer intrinsics GGML_TABLE_BEGIN(uint64_t, ksigns64, 128) 0x0000000000000000, 0xff000000000000ff, 0xff0000000000ff00, 0x000000000000ffff, 0xff00000000ff0000, 0x0000000000ff00ff, 0x0000000000ffff00, 0xff00000000ffffff, @@ -508,7 +507,6 @@ GGML_TABLE_BEGIN(uint64_t, ksigns64, 128) 0x00ffffffff000000, 0xffffffffff0000ff, 0xffffffffff00ff00, 0x00ffffffff00ffff, 0xffffffffffff0000, 0x00ffffffffff00ff, 0x00ffffffffffff00, 0xffffffffffffffff, GGML_TABLE_END() -//#endif GGML_TABLE_BEGIN(uint64_t, iq2xxs_grid, 256) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 174916bc970d7..3c766e60d6b1d 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -71,6 +71,62 @@ #define GGML_CUDA_CC_QY1 210 #define GGML_CUDA_CC_QY2 220 +#ifdef __CUDA_ARCH_LIST__ +constexpr bool ggml_cuda_has_arch_impl(int) { + return false; +} + +template +constexpr bool ggml_cuda_has_arch_impl(const int arch, const int first, Archs... rest) { + return arch == first || ggml_cuda_has_arch_impl(arch, rest...); +} + +constexpr bool ggml_cuda_has_arch(const int arch) { + return ggml_cuda_has_arch_impl(arch, __CUDA_ARCH_LIST__); +} + +static int ggml_cuda_highest_compiled_arch(const int arch) { + switch (arch) { + case 1200: if (ggml_cuda_has_arch(1200)) return 1200; [[fallthrough]]; + case 1010: if (ggml_cuda_has_arch(1010)) return 1010; [[fallthrough]]; + case 1000: if (ggml_cuda_has_arch(1000)) return 1000; [[fallthrough]]; + case 900: if (ggml_cuda_has_arch( 900)) return 900; [[fallthrough]]; + case 890: if (ggml_cuda_has_arch( 890)) return 890; [[fallthrough]]; + case 870: if (ggml_cuda_has_arch( 870)) return 870; [[fallthrough]]; + case 860: if (ggml_cuda_has_arch( 860)) return 860; [[fallthrough]]; + case 800: if (ggml_cuda_has_arch( 800)) return 800; [[fallthrough]]; + case 750: if (ggml_cuda_has_arch( 750)) return 750; [[fallthrough]]; + case 720: if (ggml_cuda_has_arch( 720)) return 720; [[fallthrough]]; + case 700: if (ggml_cuda_has_arch( 700)) return 700; [[fallthrough]]; + case 620: if (ggml_cuda_has_arch( 620)) return 620; [[fallthrough]]; + case 610: if (ggml_cuda_has_arch( 610)) return 610; [[fallthrough]]; + case 600: if (ggml_cuda_has_arch( 600)) return 600; [[fallthrough]]; + case 530: if (ggml_cuda_has_arch( 530)) return 530; [[fallthrough]]; + case 520: if (ggml_cuda_has_arch( 520)) return 520; [[fallthrough]]; + case 500: if (ggml_cuda_has_arch( 500)) return 500; [[fallthrough]]; + case 370: if (ggml_cuda_has_arch( 370)) return 370; [[fallthrough]]; + case 350: if (ggml_cuda_has_arch( 350)) return 350; [[fallthrough]]; + case 320: if (ggml_cuda_has_arch( 320)) return 320; [[fallthrough]]; + case 300: if (ggml_cuda_has_arch( 300)) return 300; [[fallthrough]]; + case 210: if (ggml_cuda_has_arch( 210)) return 210; [[fallthrough]]; + case 200: if (ggml_cuda_has_arch( 200)) return 200; [[fallthrough]]; + case 130: if (ggml_cuda_has_arch( 130)) return 130; [[fallthrough]]; + case 120: if (ggml_cuda_has_arch( 120)) return 120; [[fallthrough]]; + case 110: if (ggml_cuda_has_arch( 110)) return 110; [[fallthrough]]; + case 100: if (ggml_cuda_has_arch( 100)) return 100; + GGML_ABORT("ggml was not compiled with any CUDA arch <= %d", arch); + + default: GGML_ABORT("unknown CUDA arch: %d", arch); + } +} +#else +static int ggml_cuda_highest_compiled_arch(const int arch) { + return arch; +} +#endif // __CUDA_ARCH_LIST__ + +// --------------------------------------------------------------------------------------------------------- + #define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses #if defined(_MSC_VER) @@ -162,18 +218,22 @@ typedef float2 dfloat2; #define FLASH_ATTN_AVAILABLE #endif // !(defined(GGML_USE_MUSA) && __MUSA_ARCH__ <= GGML_CUDA_CC_QY1) -static constexpr bool fast_fp16_available(const int cc) { - return cc >= GGML_CUDA_CC_PASCAL && cc != 610; +static bool fp16_available(const int cc) { + return ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_PASCAL; +} + +static bool fast_fp16_available(const int cc) { + return fp16_available(cc) && cc != 610; } // Any FP16 tensor cores are available. -static constexpr bool fp16_mma_available(const int cc) { - return cc < GGML_CUDA_CC_OFFSET_AMD && cc >= GGML_CUDA_CC_VOLTA; +static bool fp16_mma_available(const int cc) { + return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA; } // Volta technically had FP16 tensor cores but they work very differently compared to Turing and later. -static constexpr bool new_mma_available(const int cc) { - return cc < GGML_CUDA_CC_OFFSET_AMD && cc >= GGML_CUDA_CC_TURING; +static bool new_mma_available(const int cc) { + return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING; } static constexpr __device__ int ggml_cuda_get_physical_warp_size() { diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 5b0dfacefc9da..795b720d60bcb 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -599,7 +599,7 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { case GGML_TYPE_Q5_1: return dequantize_block_cuda; case GGML_TYPE_Q8_0: - if (ggml_cuda_info().devices[ggml_cuda_get_device()].cc >= GGML_CUDA_CC_PASCAL) { + if (fp16_available(ggml_cuda_info().devices[ggml_cuda_get_device()].cc)) { return dequantize_block_q8_0_f16_cuda; } return dequantize_block_cuda; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 4dbaefdbafdf4..106257f538439 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3205,8 +3205,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g if (op->src[0]->ne[0] == 256 && op->src[1]->type == GGML_TYPE_F16 && op->src[2]->type == GGML_TYPE_F16) { return true; } - const int cc = ggml_cuda_info().devices[dev_ctx->device].cc; - return cc >= GGML_CUDA_CC_VOLTA && cc < GGML_CUDA_CC_OFFSET_AMD && op->src[1]->type == GGML_TYPE_F16 && op->src[2]->type == GGML_TYPE_F16; + return fp16_mma_available(ggml_cuda_info().devices[dev_ctx->device].cc); } case GGML_OP_CROSS_ENTROPY_LOSS: case GGML_OP_CROSS_ENTROPY_LOSS_BACK: diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 45212f66c0098..5dbddf424c4b0 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -18,7 +18,7 @@ void ggml_cuda_op_mul_mat_q( const int64_t stride00 = ne00 / ggml_blck_size(src0->type); int id = ggml_cuda_get_device(); - const int compute_capability = ggml_cuda_info().devices[id].cc; + const int cc = ggml_cuda_info().devices[id].cc; // the main device has a larger memory buffer to hold the results from all GPUs // nrows_dst == nrows of the matrix that the kernel writes into @@ -27,7 +27,8 @@ void ggml_cuda_op_mul_mat_q( // The stream-k decomposition is only faster for recent NVIDIA GPUs. // Also its fixup needs to allocate a temporary buffer in the memory pool. // There are multiple parallel CUDA streams for src1_ncols != ne11 which would introduce a race condition for this buffer. - const bool use_stream_k = compute_capability >= GGML_CUDA_CC_VOLTA && compute_capability < GGML_CUDA_CC_OFFSET_AMD && src1_ncols == ne11; + const bool use_stream_k = ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA && + cc < GGML_CUDA_CC_OFFSET_AMD && src1_ncols == ne11; const mmq_args args = {src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stride00, src1_padded_row_size, src1_ncols, ne11, nrows_dst, use_stream_k}; switch (src0->type) { @@ -136,7 +137,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) { return true; } - if (cc < GGML_CUDA_CC_DP4A) { + if (ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_DP4A) { return false; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 7a2c4d85b799f..5391542086cfd 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -86,12 +86,13 @@ struct tile_x_sizes { int sc; }; -static constexpr int get_mmq_x_max_host(const int cc) { +static int get_mmq_x_max_host(const int cc) { return new_mma_available(cc) ? 128 : + ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA && cc < GGML_CUDA_CC_OFFSET_AMD ? #ifdef GGML_CUDA_FORCE_MMQ - cc >= GGML_CUDA_CC_VOLTA && cc < GGML_CUDA_CC_OFFSET_AMD ? 128 : 64; + 128 : 64; #else - cc >= GGML_CUDA_CC_VOLTA && cc < GGML_CUDA_CC_OFFSET_AMD ? MMQ_DP4A_MAX_BATCH_SIZE : 64; + MMQ_DP4A_MAX_BATCH_SIZE : 64; #endif // GGML_CUDA_FORCE_MMQ } @@ -119,8 +120,9 @@ static constexpr __device__ int get_mmq_x_max_device() { #endif // NEW_MMA_AVAILABLE } -static constexpr int get_mmq_y_host(const int cc) { - return cc >= GGML_CUDA_CC_OFFSET_AMD ? (GGML_CUDA_CC_IS_RDNA1(cc) ? 64 : 128) : (cc >= GGML_CUDA_CC_VOLTA ? 128 : 64); +static int get_mmq_y_host(const int cc) { + return cc >= GGML_CUDA_CC_OFFSET_AMD ? (GGML_CUDA_CC_IS_RDNA1(cc) ? 64 : 128) : + (ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA ? 128 : 64); } static constexpr __device__ int get_mmq_y_device() { @@ -2828,7 +2830,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda const int mmq_x_max = get_mmq_x_max_host(cc); const int mmq_y = get_mmq_y_host(cc); const int block_num_y = (args.ne01 + mmq_y - 1) / mmq_y; - const bool use_stream_k = cc >= GGML_CUDA_CC_VOLTA && cc < GGML_CUDA_CC_OFFSET_AMD; + const bool use_stream_k = ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA && cc < GGML_CUDA_CC_OFFSET_AMD; int mmq_x_best = 0; int nparts_best = INT_MAX; From d7e36b7b50c374703e0ef9591d4d26881292d220 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Mon, 10 Feb 2025 18:29:42 +0100 Subject: [PATCH 2/4] add xy_hardware_available --- ggml/src/ggml-cuda/common.cuh | 12 +++++++++++- ggml/src/ggml-cuda/ggml-cuda.cu | 8 ++++---- ggml/src/ggml-cuda/mmq.cu | 2 +- 3 files changed, 16 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 3c766e60d6b1d..d156ac828b8bf 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -226,11 +226,21 @@ static bool fast_fp16_available(const int cc) { return fp16_available(cc) && cc != 610; } -// Any FP16 tensor cores are available. +// To be used for feature selection of external libraries, e.g. cuBLAS. +static bool fast_fp16_hardware_available(const int cc) { + return cc >= GGML_CUDA_CC_PASCAL && cc != 610; +} + +// Any FP16 tensor core instructions are available for ggml code. static bool fp16_mma_available(const int cc) { return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA; } +// To be used for feature selection of external libraries, e.g. cuBLAS. +static bool fp16_mma_hardware_available(const int cc) { + return cc < GGML_CUDA_CC_OFFSET_AMD && cc >= GGML_CUDA_CC_VOLTA; +} + // Volta technically had FP16 tensor cores but they work very differently compared to Turing and later. static bool new_mma_available(const int cc) { return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_TURING; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 106257f538439..57726a168788c 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1867,14 +1867,14 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor const int cc = ggml_cuda_info().devices[id].cc; use_mul_mat_q = use_mul_mat_q && ggml_cuda_should_use_mmq(src0->type, cc, src1->ne[1]); - any_gpus_with_slow_fp16 = any_gpus_with_slow_fp16 || !fast_fp16_available(cc); - any_gpus_without_fp16_mma = any_gpus_without_fp16_mma || !fp16_mma_available(cc); + any_gpus_with_slow_fp16 = any_gpus_with_slow_fp16 || !fast_fp16_hardware_available(cc); + any_gpus_without_fp16_mma = any_gpus_without_fp16_mma || !fp16_mma_hardware_available(cc); } } else { const int cc = ggml_cuda_info().devices[ctx.device].cc; use_mul_mat_q = use_mul_mat_q && ggml_cuda_should_use_mmq(src0->type, cc, src1->ne[1]); - any_gpus_with_slow_fp16 = any_gpus_with_slow_fp16 || !fast_fp16_available(cc); - any_gpus_without_fp16_mma = any_gpus_without_fp16_mma || !fp16_mma_available(cc); + any_gpus_with_slow_fp16 = any_gpus_with_slow_fp16 || !fast_fp16_hardware_available(cc); + any_gpus_without_fp16_mma = any_gpus_without_fp16_mma || !fp16_mma_hardware_available(cc); } // debug helpers diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 5dbddf424c4b0..5dacd131ed55c 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -146,7 +146,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) { #endif //GGML_CUDA_FORCE_MMQ if (cc < GGML_CUDA_CC_OFFSET_AMD) { - return cc < GGML_CUDA_CC_VOLTA || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; + return !fp16_mma_hardware_available(cc) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; } return (!GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc) && !GGML_CUDA_CC_IS_GCN(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; From 33c94caf13d55eb7782617c1b0110d1058b0a262 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Mon, 10 Feb 2025 18:39:05 +0100 Subject: [PATCH 3/4] revert supports_op change --- ggml/src/ggml-cuda/ggml-cuda.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 57726a168788c..c95728b08bfe8 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -3205,7 +3205,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g if (op->src[0]->ne[0] == 256 && op->src[1]->type == GGML_TYPE_F16 && op->src[2]->type == GGML_TYPE_F16) { return true; } - return fp16_mma_available(ggml_cuda_info().devices[dev_ctx->device].cc); + return fp16_mma_available(ggml_cuda_info().devices[dev_ctx->device].cc) && + op->src[1]->type == GGML_TYPE_F16 && op->src[2]->type == GGML_TYPE_F16; } case GGML_OP_CROSS_ENTROPY_LOSS: case GGML_OP_CROSS_ENTROPY_LOSS_BACK: From 03ac153a55f2996755b011d5522422204fab233f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Johannes=20G=C3=A4=C3=9Fler?= Date: Mon, 10 Feb 2025 21:33:18 +0100 Subject: [PATCH 4/4] Update ggml/src/ggml-cuda/common.cuh Co-authored-by: Diego Devesa --- ggml/src/ggml-cuda/common.cuh | 45 ++++++++++++----------------------- 1 file changed, 15 insertions(+), 30 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index d156ac828b8bf..2a3244428521b 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -85,40 +85,25 @@ constexpr bool ggml_cuda_has_arch(const int arch) { return ggml_cuda_has_arch_impl(arch, __CUDA_ARCH_LIST__); } -static int ggml_cuda_highest_compiled_arch(const int arch) { - switch (arch) { - case 1200: if (ggml_cuda_has_arch(1200)) return 1200; [[fallthrough]]; - case 1010: if (ggml_cuda_has_arch(1010)) return 1010; [[fallthrough]]; - case 1000: if (ggml_cuda_has_arch(1000)) return 1000; [[fallthrough]]; - case 900: if (ggml_cuda_has_arch( 900)) return 900; [[fallthrough]]; - case 890: if (ggml_cuda_has_arch( 890)) return 890; [[fallthrough]]; - case 870: if (ggml_cuda_has_arch( 870)) return 870; [[fallthrough]]; - case 860: if (ggml_cuda_has_arch( 860)) return 860; [[fallthrough]]; - case 800: if (ggml_cuda_has_arch( 800)) return 800; [[fallthrough]]; - case 750: if (ggml_cuda_has_arch( 750)) return 750; [[fallthrough]]; - case 720: if (ggml_cuda_has_arch( 720)) return 720; [[fallthrough]]; - case 700: if (ggml_cuda_has_arch( 700)) return 700; [[fallthrough]]; - case 620: if (ggml_cuda_has_arch( 620)) return 620; [[fallthrough]]; - case 610: if (ggml_cuda_has_arch( 610)) return 610; [[fallthrough]]; - case 600: if (ggml_cuda_has_arch( 600)) return 600; [[fallthrough]]; - case 530: if (ggml_cuda_has_arch( 530)) return 530; [[fallthrough]]; - case 520: if (ggml_cuda_has_arch( 520)) return 520; [[fallthrough]]; - case 500: if (ggml_cuda_has_arch( 500)) return 500; [[fallthrough]]; - case 370: if (ggml_cuda_has_arch( 370)) return 370; [[fallthrough]]; - case 350: if (ggml_cuda_has_arch( 350)) return 350; [[fallthrough]]; - case 320: if (ggml_cuda_has_arch( 320)) return 320; [[fallthrough]]; - case 300: if (ggml_cuda_has_arch( 300)) return 300; [[fallthrough]]; - case 210: if (ggml_cuda_has_arch( 210)) return 210; [[fallthrough]]; - case 200: if (ggml_cuda_has_arch( 200)) return 200; [[fallthrough]]; - case 130: if (ggml_cuda_has_arch( 130)) return 130; [[fallthrough]]; - case 120: if (ggml_cuda_has_arch( 120)) return 120; [[fallthrough]]; - case 110: if (ggml_cuda_has_arch( 110)) return 110; [[fallthrough]]; - case 100: if (ggml_cuda_has_arch( 100)) return 100; +constexpr int ggml_cuda_highest_compiled_arch_impl(const int arch, const int cur) { + if (cur == 0) { GGML_ABORT("ggml was not compiled with any CUDA arch <= %d", arch); + } + return cur; +} - default: GGML_ABORT("unknown CUDA arch: %d", arch); +template +constexpr int ggml_cuda_highest_compiled_arch_impl(const int arch, const int cur, const int first, Archs... rest) { + if (first <= arch && first > cur) { + return ggml_cuda_highest_compiled_arch_impl(arch, first, rest...); + } else { + return ggml_cuda_highest_compiled_arch_impl(arch, cur, rest...); } } + +constexpr int ggml_cuda_highest_compiled_arch(const int arch) { + return ggml_cuda_highest_compiled_arch_impl(arch, 0, __CUDA_ARCH_LIST__); +} #else static int ggml_cuda_highest_compiled_arch(const int arch) { return arch;