From 834f2e301ca76239109cfe8c67d13dc3b8b3ed19 Mon Sep 17 00:00:00 2001 From: Jan Morlock Date: Thu, 26 Sep 2024 23:31:05 +0200 Subject: [PATCH] change the type of u to prevent illegal memory access --- implicit/gpu/als.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/implicit/gpu/als.cu b/implicit/gpu/als.cu index 58ad047..cc2a039 100644 --- a/implicit/gpu/als.cu +++ b/implicit/gpu/als.cu @@ -35,7 +35,7 @@ least_squares_cg_kernel(int factors, size_t user_count, size_t item_count, T *X, // Stride over users in the grid: // https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/ - for (int u = blockIdx.x; u < user_count; u += gridDim.x) { + for (size_t u = blockIdx.x; u < user_count; u += gridDim.x) { T *x = &X[u * factors]; float x_value = convert(x[threadIdx.x]); @@ -102,7 +102,7 @@ least_squares_cg_kernel(int factors, size_t user_count, size_t item_count, T *X, // complain and don't let it perpetuate if (isnan(rsold)) { if (threadIdx.x == 0) { - printf("Warning NaN Detected in row %i of %lu\n", u, user_count); + printf("Warning NaN Detected in row %zu of %lu\n", u, user_count); } x_value = 0; }