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

metal : minor (readibility)

Browse source
This commit is contained in:
Georgi Gerganov 2023-09-11 10:29:34 +03:00
parent 0c17b08cc7
commit 211d82a800
No known key found for this signature in database
GPG key ID: 449E073F9DC10735
1 changed files with 3 additions and 5 deletions
Download patch file Download diff file Expand all files Collapse all files

8
ggml-metal.metal
View file

@ -213,10 +213,9 @@ kernel void kernel_diag_mask_inf_8(
dst[i+0] = src0[i+0]; dst[i+0] = src0[i+0];
dst[i+1] = src0[i+1]; dst[i+1] = src0[i+1];
int64_t i4 = 4*i; int64_t i4 = 4*i;
const int64_t i02 = i4/(ne00*ne01); const int64_t i02 = i4/(ne00*ne01); i4 -= i02*ne00*ne01;
i4 -= i02*ne00*ne01; const int64_t i01 = i4/(ne00); i4 -= i01*ne00;
const int64_t i01 = i4/ne00; const int64_t i00 = i4;
const int64_t i00 = i4 - i01*ne00;
for (int k = 3; k >= 0; --k) { for (int k = 3; k >= 0; --k) {
if (i00 + 4 + k <= n_past + i01) { if (i00 + 4 + k <= n_past + i01) {
break; break;
@ -678,7 +677,6 @@ kernel void kernel_mul_mat_f16_f32_l4(
device const half4 * x4 = (device const half4 *) (src0 + r0*nb01 + im/(ne12/ne02)*nb02); device const half4 * x4 = (device const half4 *) (src0 + r0*nb01 + im/(ne12/ne02)*nb02);
for (int r1 = 0; r1 < nrows; ++r1) { for (int r1 = 0; r1 < nrows; ++r1) {
device const float4 * y4 = (device const float4 *) (src1 + r1*nb11 + im*nb12); device const float4 * y4 = (device const float4 *) (src1 + r1*nb11 + im*nb12);
float sumf = 0; float sumf = 0;