diff --git a/ggml/src/ggml-cuda/solve_tri.cu b/ggml/src/ggml-cuda/solve_tri.cu index 726b08b99..a955d535c 100644 --- a/ggml/src/ggml-cuda/solve_tri.cu +++ b/ggml/src/ggml-cuda/solve_tri.cu @@ -44,6 +44,7 @@ static __global__ void setup_trsm_batch_pointers( } // Latency-optimized kernel for n=64, k=64 (single-token generation) +#if 0 static __global__ void solve_tri_f32_64x64_latency( const float * __restrict__ A, const float * __restrict__ B, @@ -54,8 +55,7 @@ static __global__ void solve_tri_f32_64x64_latency( const size_t nb12, const size_t nb13, const size_t nb2, - const size_t nb3) -{ + const size_t nb3) { const int batch_idx = blockIdx.x; const int lane = threadIdx.x; const int warp_id = threadIdx.y; @@ -146,6 +146,7 @@ static __global__ void solve_tri_f32_64x64_latency( X_batch[i] = sX[row * 65 + col]; } } +#endif static __global__ void solve_tri_f32_64x64_opt(const float * __restrict__ A, const float * __restrict__ B, @@ -248,6 +249,7 @@ static __global__ void solve_tri_f32_64x64_opt(const float * __restrict__ A, } } +#if 0 static __global__ void solve_tri_f32_128x128_opt(const float * __restrict__ A, const float * __restrict__ B, float * __restrict__ X, @@ -360,6 +362,7 @@ static __global__ void solve_tri_f32_128x128_opt(const float * __restrict__ A, X_batch[row * k + col] = sXt[col * 129 + row]; } } +#endif static __global__ void solve_tri_f32_256x256_tiled(const float * __restrict__ A, const float * __restrict__ B,