diff --git a/ggml-cuda.cu b/ggml-cuda.cu index 297fdbe13..67c04a176 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -1967,13 +1967,13 @@ struct mmid_row_mapping { static __global__ void k_copy_src1_to_contiguous(const char * __restrict__ src1_original, char * __restrict__ src1_contiguous, int * __restrict__ cur_src1_row, mmid_row_mapping * __restrict__ row_mapping, - const char * ids_dev, int64_t i02, size_t ids_nb1, size_t ids_nb0, + const char * __restrict ids, int64_t i02, size_t ids_nb1, size_t ids_nb0, int64_t ne11, int64_t ne10, size_t nb11, size_t nb12) { int32_t iid1 = blockIdx.x; int32_t id = blockIdx.y; - const int32_t row_id_i = *(const int32_t *) (ids_dev + iid1*ids_nb1 + id*ids_nb0); + const int32_t row_id_i = *(const int32_t *) (ids + iid1*ids_nb1 + id*ids_nb0); if (row_id_i != i02) { return; @@ -2120,7 +2120,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * CUDA_CHECK(cudaMemsetAsync(dev_cur_src1_row.get(), 0, sizeof(int), stream)); { - dim3 block_dims(std::min((uint)ne10, 512u)); + dim3 block_dims(std::min((unsigned int)ne10, 512u)); dim3 grid_dims(ids->ne[1], n_ids); k_copy_src1_to_contiguous<<>>( src1_original, src1_contiguous.get(), @@ -2151,7 +2151,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * #ifndef MMID_MEMCPY { - dim3 block_dims(std::min((uint)ne0, 512u)); + dim3 block_dims(std::min((unsigned int)ne0, 512u)); dim3 grid_dims(num_src1_rows); k_copy_dst_from_contiguous<<>>( dst_original, dst_contiguous.get(),