root/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2024-12-24 10:24:35 +00:00
Code Issues Actions 5 Packages Projects Releases Wiki Activity

cuda : fix missing const qualifier in casts (#2027)

Browse Source
This commit is contained in:
Salvador E. Tropea 2023-06-28 14:26:26 -03:00 committed by GitHub
parent b922bc351b
commit 6432aabb6d
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 4 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
ggml-cuda.cu
View File

@ -1244,7 +1244,7 @@ static __global__ void dequantize_mul_mat_vec(const void * vx, const dfloat * y,
} }
static __global__ void mul_mat_p021_f16_f32(const void * vx, const float * y, float * dst, const int ncols_x, const int nrows_x, const int nchannels_x) { static __global__ void mul_mat_p021_f16_f32(const void * vx, const float * y, float * dst, const int ncols_x, const int nrows_x, const int nchannels_x) {
const half * x = (half *) vx; const half * x = (const half *) vx;
const int row_x = blockDim.y*blockIdx.y + threadIdx.y; const int row_x = blockDim.y*blockIdx.y + threadIdx.y;
const int channel = blockDim.z*blockIdx.z + threadIdx.z; const int channel = blockDim.z*blockIdx.z + threadIdx.z;
@ -1294,7 +1294,7 @@ static __global__ void mul_mat_vec_nc_f16_f32( // nc == non-contiguous
const void * vx, const float * y, float * dst, const int ncols_x, const int nrows_x, const void * vx, const float * y, float * dst, const int ncols_x, const int nrows_x,
const int row_stride_x, const int nchannels_x, const int channel_stride_x) { const int row_stride_x, const int nchannels_x, const int channel_stride_x) {
const half * x = (half *) vx; const half * x = (const half *) vx;
const int row_x = blockDim.y*blockIdx.y + threadIdx.y; const int row_x = blockDim.y*blockIdx.y + threadIdx.y;
const int channel = blockDim.z*blockIdx.z + threadIdx.z; const int channel = blockDim.z*blockIdx.z + threadIdx.z;
@ -1337,14 +1337,14 @@ static __global__ void mul_mat_vec_nc_f16_f32( // nc == non-contiguous
} }
static __device__ void cpy_1_f32_f32(const char * cxi, char * cdsti) { static __device__ void cpy_1_f32_f32(const char * cxi, char * cdsti) {
const float * xi = (float *) cxi; const float * xi = (const float *) cxi;
float * dsti = (float *) cdsti; float * dsti = (float *) cdsti;
*dsti = *xi; *dsti = *xi;
} }
static __device__ void cpy_1_f32_f16(const char * cxi, char * cdsti) { static __device__ void cpy_1_f32_f16(const char * cxi, char * cdsti) {
const float * xi = (float *) cxi; const float * xi = (const float *) cxi;
half * dsti = (half *) cdsti; half * dsti = (half *) cdsti;
*dsti = __float2half(*xi); *dsti = __float2half(*xi);