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

Fix boundary check in conv2d_implicit_kernel to include channel limits

This commit is contained in:
bssrdf 2025-09-04 15:32:29 -04:00
parent 6d84cbb5ab
commit 5ffe97be9c
1 changed files with 2 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -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];
}