Fix boundary check in conv2d_implicit_kernel to include channel limits

bssrdf · bssrdf · commit 5ffe97be9c35 · 2025-09-04T15:32:29.000-04:00
diff --git 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];
             }

Original file line number	Diff line number	Diff line change
`@@ -116,7 +116,7 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input,`
`116`	`116`	`int curH = posh_ori[i] + curR * param.d_h; // input h`
`117`	`117`	`int curW = posw_ori[i] + curS * param.d_w; // input w`
`118`	`118`	`int inOffsetTmp = curC * inChannelOffset + curH * param.w + curW;`
`119`		`- if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h)`
	`119`	`+ if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h && curC < param.c)`
`120`	`120`	`{`
`121`	`121`	`input_ldg_reg[i] = input[inOffset + inOffsetTmp];`
`122`	`122`	`}`
`@@ -175,7 +175,7 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input,`
`175`	`175`	`int curH = posh_ori[i] + curR * param.d_h; // input h`
`176`	`176`	`int curW = posw_ori[i] + curS * param.d_w; // input w`
`177`	`177`	`int inOffsetTmp = curC * inChannelOffset + curH * param.w + curW;`
`178`		`- if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h)`
	`178`	`+ if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h && curC < param.c)`
`179`	`179`	`{`
`180`	`180`	`input_ldg_reg[i] = input[inOffset + inOffsetTmp];`
`181`	`181`	`}`