From 5ffe97be9c35169aea1e451426eb53e5430f7d24 Mon Sep 17 00:00:00 2001 From: bssrdf Date: Thu, 4 Sep 2025 15:32:29 -0400 Subject: [PATCH] Fix boundary check in conv2d_implicit_kernel to include channel limits --- ggml/src/ggml-cuda/conv2d-implicit.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/conv2d-implicit.cu b/ggml/src/ggml-cuda/conv2d-implicit.cu index 4f452ab98b..d9fabd9657 100644 --- a/ggml/src/ggml-cuda/conv2d-implicit.cu +++ b/ggml/src/ggml-cuda/conv2d-implicit.cu @@ -116,7 +116,7 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input, int curH = posh_ori[i] + curR * param.d_h; // input h int curW = posw_ori[i] + curS * param.d_w; // input w int inOffsetTmp = curC * inChannelOffset + curH * param.w + curW; - if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h) + if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h && curC < param.c) { input_ldg_reg[i] = input[inOffset + inOffsetTmp]; } @@ -175,7 +175,7 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input, int curH = posh_ori[i] + curR * param.d_h; // input h int curW = posw_ori[i] + curS * param.d_w; // input w int inOffsetTmp = curC * inChannelOffset + curH * param.w + curW; - if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h) + if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h && curC < param.c) { input_ldg_reg[i] = input[inOffset + inOffsetTmp]; }