summaryrefslogtreecommitdiff
path: root/ggml-cuda.cu
diff options
context:
space:
mode:
authorGeorgi Gerganov <ggerganov@gmail.com>2023-12-22 17:53:43 +0200
committerGitHub <noreply@github.com>2023-12-22 17:53:43 +0200
commitba661751322a7c201fd3bef71af077c5aebfaa2a (patch)
treeaf645f9d0eefb71ace31501540ab4263ecc6aaf8 /ggml-cuda.cu
parenta55876955b1a83464171de8d578d3ab062a7b62d (diff)
sync : ggml (fix im2col) (#4591)
* cuda : fix im2col_f32_f16 (ggml/#658) ggml-ci * ggml-alloc : fix ggml_tallocr_is_own --------- Co-authored-by: leejet <leejet714@gmail.com>
Diffstat (limited to 'ggml-cuda.cu')
-rw-r--r--ggml-cuda.cu8
1 files changed, 4 insertions, 4 deletions
diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index b124774a..7c2a834e 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -5273,17 +5273,17 @@ static __global__ void im2col_f32_f16(
const int ky = (i - kd) / OW;
const int ix = i % OW;
- const int iiw = ix * s0 + kx * d0 - p0;
- const int iih = blockIdx.y * s1 + ky * d1 - p1;
+ const int64_t iiw = ix * s0 + kx * d0 - p0;
+ const int64_t iih = blockIdx.y * s1 + ky * d1 - p1;
- const int offset_dst =
+ const int64_t offset_dst =
(blockIdx.y * OW + ix) * CHW +
(blockIdx.z * (KW * KH) + ky * KW + kx);
if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW) {
dst[offset_dst] = __float2half(0.0f);
} else {
- const int offset_src = blockIdx.z * offset_delta;
+ const int64_t offset_src = blockIdx.z * offset_delta;
dst[offset_dst] = __float2half(x[offset_src + iih * IW + iiw]);
}
}