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

remove compile flag

Browse source
This commit is contained in:
Georgi Gerganov 2024-11-07 19:18:31 +02:00
parent a6c8dbfa5d
commit 022e5e90e9
No known key found for this signature in database
GPG key ID: 449E073F9DC10735
2 changed files with 5 additions and 12 deletions
Download patch file Download diff file Expand all files Collapse all files

9
ggml/src/ggml-metal.m
View file

@ -504,10 +504,6 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
// dictionary of preprocessor macros // dictionary of preprocessor macros
NSMutableDictionary * prep = [NSMutableDictionary dictionary]; NSMutableDictionary * prep = [NSMutableDictionary dictionary];
#if defined(GGML_METAL_FORCE_FATTN_PREC_F16)
[prep setObject:@"1" forKey:@"GGML_METAL_FORCE_FATTN_PREC_F16"];
#endif
MTLCompileOptions * options = [MTLCompileOptions new]; MTLCompileOptions * options = [MTLCompileOptions new];
options.preprocessorMacros = prep; options.preprocessorMacros = prep;
@ -558,11 +554,6 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
} }
} }
#if defined(GGML_METAL_FORCE_FATTN_PREC_F16)
GGML_LOG_INFO("%s: GGML_METAL_FORCE_FATTN_PREC_F16 = yes\n", __func__);
#else
GGML_LOG_INFO("%s: GGML_METAL_FORCE_FATTN_PREC_F16 = no\n", __func__);
#endif
GGML_LOG_INFO("%s: simdgroup reduction = %s\n", __func__, ctx_dev->has_simdgroup_reduction ? "true" : "false"); GGML_LOG_INFO("%s: simdgroup reduction = %s\n", __func__, ctx_dev->has_simdgroup_reduction ? "true" : "false");
GGML_LOG_INFO("%s: simdgroup matrix mul. = %s\n", __func__, ctx_dev->has_simdgroup_mm ? "true" : "false"); GGML_LOG_INFO("%s: simdgroup matrix mul. = %s\n", __func__, ctx_dev->has_simdgroup_mm ? "true" : "false");
GGML_LOG_INFO("%s: bfloat = %s\n", __func__, ctx_dev->has_bfloat ? "true" : "false"); GGML_LOG_INFO("%s: bfloat = %s\n", __func__, ctx_dev->has_bfloat ? "true" : "false");

8
ggml/src/ggml-metal.metal
View file

@ -3021,9 +3021,11 @@ kernel void kernel_flash_attn_ext(
} }
// cast qk_t -> s_t // cast qk_t -> s_t
s8x8_t mqks(1.0f); //s8x8_t mqks(1.0f);
simdgroup_multiply(mqks, mqk, mqks); //simdgroup_multiply(mqks, mqk, mqks);
simdgroup_store(mqks, ss + 8*cc, TS, 0, false); //simdgroup_store(mqks, ss + 8*cc, TS, 0, false);
simdgroup_store(mqk, ss + 8*cc, TS, 0, false);
} }
} }