vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

fix mixtral

Browse source
This commit is contained in:
slaren 2023-12-25 13:24:36 +01:00
parent 6f35a4a6e9
commit 1659cd1bc4
1 changed files with 12 additions and 2 deletions
Download patch file Download diff file Expand all files Collapse all files

10
ggml-cuda.cu
View file

@ -8258,8 +8258,18 @@ static void ggml_cuda_op_mul_mat(
float * dhf_dst_i = (float *) ((char *) dst_off_device + i02*nb2 + i03*nb3); float * dhf_dst_i = (float *) ((char *) dst_off_device + i02*nb2 + i03*nb3);
GGML_ASSERT(dst->nb[1] == ne0*sizeof(float)); GGML_ASSERT(dst->nb[1] == ne0*sizeof(float));
dhf_dst_i += src1_col_0*ne0 + row_low[id]; dhf_dst_i += src1_col_0*ne0 + row_low[id];
if (kind == cudaMemcpyDeviceToDevice && id != g_main_device) {
// there is no cudaMemcpy2DPeerAsync so we need to copy each row separately
for (int64_t i = 0; i < src1_ncols; ++i) {
CUDA_CHECK(cudaMemcpyPeerAsync(dhf_dst_i + i*ne0, g_main_device,
dst_dd_i + i*row_diff, id,
row_diff*sizeof(float), stream));
}
} else {
CUDA_CHECK(cudaMemcpy2DAsync(dhf_dst_i, ne0*sizeof(float), dst_dd_i, row_diff*sizeof(float), CUDA_CHECK(cudaMemcpy2DAsync(dhf_dst_i, ne0*sizeof(float), dst_dd_i, row_diff*sizeof(float),
row_diff*sizeof(float), src1_ncols, kind, stream)); row_diff*sizeof(float), src1_ncols, kind, stream));
}
} else { } else {
float * dhf_dst_i = (float *) ((char *) dst_off_device + i02*nb2 + i03*nb3); float * dhf_dst_i = (float *) ((char *) dst_off_device + i02*nb2 + i03*nb3);
GGML_ASSERT(dst->nb[1] == ne0*sizeof(float)); GGML_ASSERT(dst->nb[1] == ne0*sizeof(float));