happyz
/
llama.cpp
mirror of https://github.com/ggerganov/llama.cpp.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

fix to work for Turing

This commit is contained in:
bssrdf 2026-01-29 16:08:56 -05:00
parent 805e9ac6a8
commit 75e5a9bd01
1 changed files with 3 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
ggml/src/ggml-cuda/conv2d-implicit.cu
View File

@ -344,7 +344,7 @@ __device__ __forceinline__ void ldmatrix_a(
const half* src,
half (&reg)[mma_tiles_per_warp_m][mma_tiles_per_warp_k][4]
){
#if __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
#ifdef CP_ASYNC_AVAILABLE
static_assert(mma_tiles_per_warp_m == 8, "mma_tiles_per_warp_m must be 8");
static_assert(mma_tiles_per_warp_k == 4, "mma_tiles_per_warp_k must be 4");
@ -503,7 +503,7 @@ __device__ __forceinline__ void ldmatrix_b(
const half* src,
half (&reg)[mma_tiles_per_warp_k][mma_tiles_per_warp_n][2]
){
#if __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
#ifdef CP_ASYNC_AVAILABLE
static_assert(mma_tiles_per_warp_k == 4, "mma_tiles_per_warp_k must be 4");
static_assert(mma_tiles_per_warp_n == 8, "mma_tiles_per_warp_n must be 8");
@ -981,8 +981,7 @@ static void launch_conv2d_implicit_split_kernel(ggml_backend_cuda_context & ctx,
static void conv2d_implicit_cuda_f16(ggml_backend_cuda_context & ctx, const float * X_D, const half * K_D, float * Y_D, int cc, param_t P, cudaStream_t st) {
// if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && P.c % 8 == 0 && (P.r > 1 || P.s > 1)) {
if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && P.c % 8 == 0 && (P.r <= 32 && P.s <= 32)) {
if (GGML_CUDA_CC_IS_NVIDIA(cc) && ampere_mma_available(cc) && P.c % 8 == 0 && (P.r <= 32 && P.s <= 32)) {
int id = ggml_cuda_get_device();