diff --git a/Makefile b/Makefile index 5bc806bd1abf..ae239b581b42 100644 --- a/Makefile +++ b/Makefile @@ -139,10 +139,10 @@ NOAVX2_BUILD = OPENBLAS_NOAVX2_BUILD = ifeq ($(OS),Windows_NT) - OPENBLAS_BUILD = $(CXX) $(CXXFLAGS) ggml_openblas.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o lib/libopenblas.lib -shared -o koboldcpp_openblas.dll $(LDFLAGS) - CLBLAST_BUILD = $(CXX) $(CXXFLAGS) ggml_clblast.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o lib/OpenCL.lib lib/clblast.lib -shared -o koboldcpp_clblast.dll $(LDFLAGS) - OPENBLAS_NOAVX2_BUILD = $(CXX) $(CXXFLAGS) ggml_openblas_noavx2.o ggml_v1_noavx2.o expose.o common.o llama_adapter.o gpttype_adapter.o lib/libopenblas.lib -shared -o koboldcpp_openblas_noavx2.dll $(LDFLAGS) - NOAVX2_BUILD = $(CXX) $(CXXFLAGS) ggml_noavx2.o ggml_v1_noavx2.o expose.o common.o llama_adapter.o gpttype_adapter.o -shared -o koboldcpp_noavx2.dll $(LDFLAGS) + OPENBLAS_BUILD = $(CXX) $(CXXFLAGS) $^ lib/libopenblas.lib -shared -o $@ $(LDFLAGS) + CLBLAST_BUILD = $(CXX) $(CXXFLAGS) $^ lib/OpenCL.lib lib/clblast.lib -shared -o $@ $(LDFLAGS) + OPENBLAS_NOAVX2_BUILD = $(CXX) $(CXXFLAGS) $^ lib/libopenblas.lib -shared -o $@ $(LDFLAGS) + NOAVX2_BUILD = $(CXX) $(CXXFLAGS) $^ -shared -o $@ $(LDFLAGS) else ifndef LLAMA_OPENBLAS ifndef LLAMA_CLBLAST @@ -166,101 +166,101 @@ $(info I CC: $(CCV)) $(info I CXX: $(CXXV)) $(info ) -default: llamalib llamalib_noavx2 llamalib_openblas llamalib_openblas_noavx2 llamalib_clblast -simple: llamalib llamalib_noavx2 -dev: llamalib_openblas +default: koboldcpp.dll koboldcpp_noavx2.dll koboldcpp_openblas.dll koboldcpp_openblas_noavx2.dll koboldcpp_clblast.dll +simple: koboldcpp.dll koboldcpp_noavx2.dll +dev: koboldcpp_openblas.dll # # Build library # ggml.o: ggml.c ggml.h - $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -c ggml.c -o ggml.o + $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -c $< -o $@ ggml_openblas.o: ggml.c ggml.h - $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -DGGML_USE_OPENBLAS -c ggml.c -o ggml_openblas.o + $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -DGGML_USE_OPENBLAS -c $< -o $@ ggml_noavx2.o: ggml.c ggml.h - $(CC) $(CFLAGS) -c ggml.c -o ggml_noavx2.o + $(CC) $(CFLAGS) -c $< -o $@ ggml_openblas_noavx2.o: ggml.c ggml.h - $(CC) $(CFLAGS) -DGGML_USE_OPENBLAS -c ggml.c -o ggml_openblas_noavx2.o + $(CC) $(CFLAGS) -DGGML_USE_OPENBLAS -c $< -o $@ ggml_clblast.o: ggml.c ggml.h - $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -DGGML_USE_OPENBLAS -DGGML_USE_CLBLAST -c ggml.c -o ggml_clblast.o + $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -DGGML_USE_OPENBLAS -DGGML_USE_CLBLAST -c $< -o $@ ggml_v1.o: otherarch/ggml_v1.c otherarch/ggml_v1.h - $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -c otherarch/ggml_v1.c -o ggml_v1.o + $(CC) $(CFLAGS) $(BONUSCFLAGS1) $(BONUSCFLAGS2) -c $< -o $@ ggml_v1_noavx2.o: otherarch/ggml_v1.c otherarch/ggml_v1.h - $(CC) $(CFLAGS) $(BONUSCFLAGS1) -c otherarch/ggml_v1.c -o ggml_v1_noavx2.o + $(CC) $(CFLAGS) $(BONUSCFLAGS1) -c $< -o $@ llama.o: llama.cpp llama.h llama_util.h - $(CXX) $(CXXFLAGS) -c llama.cpp -o llama.o + $(CXX) $(CXXFLAGS) -c $< -o $@ common.o: examples/common.cpp examples/common.h - $(CXX) $(CXXFLAGS) -c examples/common.cpp -o common.o + $(CXX) $(CXXFLAGS) -c $< -o $@ expose.o: expose.cpp expose.h - $(CXX) $(CXXFLAGS) -c expose.cpp -o expose.o + $(CXX) $(CXXFLAGS) -c $< -o $@ -llama_adapter.o: - $(CXX) $(CXXFLAGS) -c llama_adapter.cpp -o llama_adapter.o +llama_adapter.o: llama_adapter.cpp + $(CXX) $(CXXFLAGS) -c $< -o $@ -gpttype_adapter.o: - $(CXX) $(CXXFLAGS) -c gpttype_adapter.cpp -o gpttype_adapter.o +gpttype_adapter.o: gpttype_adapter.cpp + $(CXX) $(CXXFLAGS) -c $< -o $@ clean: rm -vf *.o main quantize_llama quantize_gpt2 quantize_gptj quantize-stats perplexity embedding benchmark-q4_0-matmult main.exe quantize_llama.exe quantize_gptj.exe quantize_gpt2.exe koboldcpp.dll koboldcpp_openblas.dll koboldcpp_noavx2.dll koboldcpp_openblas_noavx2.dll koboldcpp_clblast.dll gptj.exe gpt2.exe main: examples/main/main.cpp ggml.o llama.o common.o - $(CXX) $(CXXFLAGS) examples/main/main.cpp ggml.o llama.o common.o -o main $(LDFLAGS) + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) @echo @echo '==== Run ./main -h for help. ====' @echo -llamalib: ggml.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o - $(CXX) $(CXXFLAGS) ggml.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o -shared -o koboldcpp.dll $(LDFLAGS) +koboldcpp.dll: ggml.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o + $(CXX) $(CXXFLAGS) $^ -shared -o $@ $(LDFLAGS) -llamalib_openblas: ggml_openblas.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o +koboldcpp_openblas.dll: ggml_openblas.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o $(OPENBLAS_BUILD) - -llamalib_noavx2: ggml_noavx2.o ggml_v1_noavx2.o expose.o common.o llama_adapter.o gpttype_adapter.o + +koboldcpp_noavx2.dll: ggml_noavx2.o ggml_v1_noavx2.o expose.o common.o llama_adapter.o gpttype_adapter.o $(NOAVX2_BUILD) -llamalib_openblas_noavx2: ggml_openblas_noavx2.o ggml_v1_noavx2.o expose.o common.o llama_adapter.o gpttype_adapter.o +koboldcpp_openblas_noavx2.dll: ggml_openblas_noavx2.o ggml_v1_noavx2.o expose.o common.o llama_adapter.o gpttype_adapter.o $(OPENBLAS_NOAVX2_BUILD) -llamalib_clblast: ggml_clblast.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o +koboldcpp_clblast.dll: ggml_clblast.o ggml_v1.o expose.o common.o llama_adapter.o gpttype_adapter.o $(CLBLAST_BUILD) quantize_llama: examples/quantize/quantize.cpp ggml.o llama.o - $(CXX) $(CXXFLAGS) examples/quantize/quantize.cpp ggml.o llama.o -o quantize_llama $(LDFLAGS) + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) quantize-stats: examples/quantize-stats/quantize-stats.cpp ggml.o llama.o - $(CXX) $(CXXFLAGS) examples/quantize-stats/quantize-stats.cpp ggml.o llama.o -o quantize-stats $(LDFLAGS) + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) -quantize_gptj: ggml.o llama.o - $(CXX) $(CXXFLAGS) otherarch/gptj_quantize.cpp ggml.o llama.o -o quantize_gptj $(LDFLAGS) +quantize_gptj: ggml.o llama.o otherarch/gptj_quantize.cpp + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) -quantize_gpt2: ggml.o llama.o - $(CXX) $(CXXFLAGS) otherarch/gpt2_quantize.cpp ggml.o llama.o -o quantize_gpt2 $(LDFLAGS) +quantize_gpt2: ggml.o llama.o otherarch/gpt2_quantize.cpp + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) perplexity: examples/perplexity/perplexity.cpp ggml.o llama.o common.o - $(CXX) $(CXXFLAGS) examples/perplexity/perplexity.cpp ggml.o llama.o common.o -o perplexity $(LDFLAGS) + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) embedding: examples/embedding/embedding.cpp ggml.o llama.o common.o - $(CXX) $(CXXFLAGS) examples/embedding/embedding.cpp ggml.o llama.o common.o -o embedding $(LDFLAGS) + $(CXX) $(CXXFLAGS) $^ -o $@ $(LDFLAGS) libllama.so: llama.o ggml.o - $(CXX) $(CXXFLAGS) -shared -fPIC -o libllama.so llama.o ggml.o $(LDFLAGS) + $(CXX) $(CXXFLAGS) -shared -fPIC -o $@ $^ $(LDFLAGS) # # Tests # -benchmark: ggml.o - $(CXX) $(CXXFLAGS) examples/benchmark/benchmark-q4_0-matmult.c ggml.o -o benchmark-q4_0-matmult $(LDFLAGS) +benchmark: examples/benchmark/benchmark-q4_0-matmult.c ggml.o + $(CXX) $(CXXFLAGS) $^ -o benchmark-q4_0-matmult $(LDFLAGS) ./benchmark-q4_0-matmult .PHONY: tests diff --git a/examples/quantize-stats/quantize-stats.cpp b/examples/quantize-stats/quantize-stats.cpp index c786fe208c78..0503009319e4 100644 --- a/examples/quantize-stats/quantize-stats.cpp +++ b/examples/quantize-stats/quantize-stats.cpp @@ -16,9 +16,6 @@ #include #include -static const char * type_strs[] = { "q4_0", "q4_1", "i8", "i16", "i32", "f16", "f32" }; -static_assert(sizeof(type_strs) == GGML_TYPE_COUNT * sizeof(char *), "Incomplete type list"); - struct quantize_stats_params { std::string model = "models/7B/ggml-model-f16.bin"; bool verbose = false; @@ -224,7 +221,7 @@ int main(int argc, char ** argv) { break; } int j; - for (j = 0; j < GGML_TYPE_COUNT && strcmp(argv[i], type_strs[j]) != 0; j++) { + for (j = 0; j < GGML_TYPE_COUNT && strcmp(argv[i], ggml_type_name((ggml_type) i)) != 0; j++) { // find match } if (j < GGML_TYPE_COUNT) { @@ -279,7 +276,7 @@ int main(int argc, char ** argv) { continue; } if (params.verbose) { - printf("%s: type %s, size %" PRId64 "\n", kv_tensor.first.c_str(), type_strs[kv_tensor.second->type], ggml_nelements(kv_tensor.second)); + printf("%s: type %s, size %" PRId64 "\n", kv_tensor.first.c_str(), ggml_type_name(kv_tensor.second->type), ggml_nelements(kv_tensor.second)); } if (kv_tensor.second->type == GGML_TYPE_F16) { is_f16 = true; @@ -304,13 +301,14 @@ int main(int argc, char ** argv) { // loop throught quantization types for (int i = 0; i < GGML_TYPE_COUNT; i++) { + const ggml_type type = (ggml_type) i; if (!params.include_types.empty() && std::find(params.include_types.begin(), params.include_types.end(), i) == params.include_types.end()) { continue; } quantize_fns_t qfns = ggml_internal_get_quantize_fn(i); if (qfns.quantize_row_q && qfns.dequantize_row_q) { if (params.verbose) { - printf("testing %s ...\n", type_strs[i]); + printf("testing %s ...\n", ggml_type_name(type)); } error_stats global_stats {}; @@ -322,7 +320,7 @@ int main(int argc, char ** argv) { if (params.verbose) { printf(" %s ...\n", kv_tensor.first.c_str()); } - std::string layer_name { type_strs[i] }; + std::string layer_name { ggml_type_name(type) }; layer_name += "::" + kv_tensor.first; test_roundtrip_on_layer( layer_name, @@ -337,7 +335,7 @@ int main(int argc, char ** argv) { ); } - print_error_stats(type_strs[i], global_stats, params.print_histogram); + print_error_stats(ggml_type_name(type), global_stats, params.print_histogram); } } diff --git a/expose.cpp b/expose.cpp index 893f427a0adc..d8b62b3e1b2f 100644 --- a/expose.cpp +++ b/expose.cpp @@ -23,6 +23,8 @@ extern "C" { + std::string platformenv, deviceenv; + //return val: 0=fail, 1=(original ggml, alpaca), 2=(ggmf), 3=(ggjt) static FileFormat file_format = FileFormat::BADFORMAT; @@ -33,15 +35,15 @@ extern "C" //first digit is whether configured, second is platform, third is devices int parseinfo = inputs.clblast_info; - + std::string usingclblast = "KCPP_CLBLAST_CONFIGURED="+std::to_string(parseinfo>0?1:0); putenv((char*)usingclblast.c_str()); - parseinfo = parseinfo%100; //keep last 2 digits + parseinfo = parseinfo%100; //keep last 2 digits int platform = parseinfo/10; int devices = parseinfo%10; - std::string platformenv = "KCPP_CLBLAST_PLATFORM="+std::to_string(platform); - std::string deviceenv = "KCPP_CLBLAST_DEVICES="+std::to_string(devices); + platformenv = "KCPP_CLBLAST_PLATFORM="+std::to_string(platform); + deviceenv = "KCPP_CLBLAST_DEVICES="+std::to_string(devices); putenv((char*)platformenv.c_str()); putenv((char*)deviceenv.c_str()); @@ -61,7 +63,7 @@ extern "C" printf("\n---\nRetrying as GPT-J model: (ver %d)\nAttempting to Load...\n---\n", file_format); lr = gpttype_load_model(inputs, file_format); } - + if (lr == ModelLoadResult::FAIL || lr == ModelLoadResult::RETRY_LOAD) { return false; @@ -92,14 +94,14 @@ extern "C" } else { - printf("\n---\nIdentified as LLAMA model: (ver %d)\nAttempting to Load...\n---\n", file_format); + printf("\n---\nIdentified as LLAMA model: (ver %d)\nAttempting to Load...\n---\n", file_format); return llama_load_model(inputs, file_format); } } generation_outputs generate(const generation_inputs inputs, generation_outputs &output) { - if (file_format == FileFormat::GPTJ_1 || file_format == FileFormat::GPTJ_2 || file_format==FileFormat::GPTJ_3 + if (file_format == FileFormat::GPTJ_1 || file_format == FileFormat::GPTJ_2 || file_format==FileFormat::GPTJ_3 || file_format==FileFormat::GPT2_1 || file_format==FileFormat::GPT2_2 ) { return gpttype_generate(inputs, output); @@ -107,6 +109,6 @@ extern "C" else { return llama_generate(inputs, output); - } + } } } diff --git a/expose.h b/expose.h index 66dfde66afb8..67f1e1937386 100644 --- a/expose.h +++ b/expose.h @@ -11,6 +11,7 @@ struct load_model_inputs const bool use_mmap; const bool use_smartcontext; const int clblast_info = 0; + const int blasbatchsize = 512; }; struct generation_inputs { diff --git a/ggml.c b/ggml.c index 6f67be22f452..7b00b6ad3afe 100644 --- a/ggml.c +++ b/ggml.c @@ -2671,6 +2671,18 @@ static const size_t GGML_TYPE_SIZE[GGML_TYPE_COUNT] = { }; static_assert(GGML_TYPE_COUNT == 7, "GGML_TYPE_SIZE is outdated"); + +static const char * GGML_TYPE_NAME[GGML_TYPE_COUNT] = { + [GGML_TYPE_F32] = "f32", + [GGML_TYPE_F16] = "f16", + [GGML_TYPE_Q4_0] = "q4_0", + [GGML_TYPE_Q4_1] = "q4_1", + [GGML_TYPE_I8] = "i8", + [GGML_TYPE_I16] = "i16", + [GGML_TYPE_I32] = "i32", +}; +static_assert(GGML_TYPE_COUNT == 7, "GGML_TYPE_NAME is outdated"); + static const char * GGML_OP_LABEL[GGML_OP_COUNT] = { "NONE", @@ -2712,9 +2724,12 @@ static const char * GGML_OP_LABEL[GGML_OP_COUNT] = { "FLASH_ATTN", "FLASH_FF", + + "MAP_UNARY", + "MAP_BINARY", }; -static_assert(GGML_OP_COUNT == 36, "GGML_OP_COUNT != 36"); +static_assert(GGML_OP_COUNT == 38, "GGML_OP_COUNT != 38"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -2757,9 +2772,12 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "flash_attn(x)", "flash_ff(x)", + + "f(x)", + "f(x,y)", }; -static_assert(GGML_OP_COUNT == 36, "GGML_OP_COUNT != 36"); +static_assert(GGML_OP_COUNT == 38, "GGML_OP_COUNT != 38"); static_assert(sizeof(struct ggml_object)%GGML_MEM_ALIGN == 0, "ggml_object size must be a multiple of GGML_MEM_ALIGN"); static_assert(sizeof(struct ggml_tensor)%GGML_MEM_ALIGN == 0, "ggml_tensor size must be a multiple of GGML_MEM_ALIGN"); @@ -2889,6 +2907,11 @@ float ggml_type_sizef(enum ggml_type type) { return ((float)(GGML_TYPE_SIZE[type]))/GGML_BLCK_SIZE[type]; } +const char * ggml_type_name(enum ggml_type type) { + return GGML_TYPE_NAME[type]; +} + + size_t ggml_element_size(const struct ggml_tensor * tensor) { return GGML_TYPE_SIZE[tensor->type]; } @@ -4907,6 +4930,90 @@ struct ggml_tensor * ggml_flash_ff( return result; } +// ggml_map_unary + +struct ggml_tensor * ggml_map_unary_impl_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + const ggml_unary_op_f32_t fun, + bool inplace) { + bool is_node = false; + + if (!inplace && a->grad) { + is_node = true; + } + + struct ggml_tensor * addr_tensor = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, sizeof(void *) / sizeof(int32_t)); + *((void (**)(void))addr_tensor->data) = (void (*)(void))fun; + struct ggml_tensor *result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a); + + result->op = GGML_OP_MAP_UNARY; + result->grad = is_node ? ggml_dup_tensor(ctx, result) : NULL; + result->src0 = a; + result->opt[0] = addr_tensor; + + return result; +} + +struct ggml_tensor * ggml_map_unary_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + const ggml_unary_op_f32_t fun) { + return ggml_map_unary_impl_f32(ctx, a, fun, false); +} + +struct ggml_tensor * ggml_map_unary_inplace_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + const ggml_unary_op_f32_t fun) { + return ggml_map_unary_impl_f32(ctx, a, fun, true); +} + +// ggml_map_binary + +struct ggml_tensor * ggml_map_binary_impl_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + const ggml_binary_op_f32_t fun, + bool inplace) { + GGML_ASSERT(ggml_are_same_shape(a, b)); + + bool is_node = false; + + if (!inplace && (a->grad || b->grad)) { + is_node = true; + } + + struct ggml_tensor * addr_tensor = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, sizeof(void *) / sizeof(int32_t)); + *((void (**)(void))addr_tensor->data) = (void (*)(void))fun; + struct ggml_tensor *result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a); + + result->op = GGML_OP_MAP_BINARY; + result->grad = is_node ? ggml_dup_tensor(ctx, result) : NULL; + result->src0 = a; + result->src1 = b; + result->opt[0] = addr_tensor; + + return result; +} + +struct ggml_tensor * ggml_map_binary_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + const ggml_binary_op_f32_t fun) { + return ggml_map_binary_impl_f32(ctx, a, b, fun, false); +} + +struct ggml_tensor * ggml_map_binary_inplace_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + const ggml_binary_op_f32_t fun) { + return ggml_map_binary_impl_f32(ctx, a, b, fun, true); +} + //////////////////////////////////////////////////////////////////////////////// void ggml_set_param( @@ -6528,7 +6635,8 @@ static void ggml_compute_forward_mul_mat_f32( ne11, ne01, ne10, 1.0f, y, ne10, x, ne00, - 0.0f, d, ne01); + 0.0f, d, ne01, + params->type); } } @@ -6700,7 +6808,8 @@ static void ggml_compute_forward_mul_mat_f16_f32( ne11, ne01, ne10, 1.0f, y, ne10, x, ne00, - 0.0f, d, ne01); + 0.0f, d, ne01, + params->type); } } @@ -6895,6 +7004,7 @@ static void ggml_compute_forward_mul_mat_q_f32( for (int64_t i03 = 0; i03 < ne03; i03++) { for (int64_t i02 = 0; i02 < ne02; i02++) { +#ifndef GGML_USE_CLBLAST { size_t id = 0; for (int64_t i01 = 0; i01 < ne01; ++i01) { @@ -6904,6 +7014,9 @@ static void ggml_compute_forward_mul_mat_q_f32( } const float * x = wdata; +#else + const void* x = src0->data + i03*nb03 + i02*nb02; +#endif const float * y = (float *) ((char *) src1->data + i02*nb12 + i03*nb13); float * d = (float *) ((char *) dst->data + i02*nb2 + i03*nb3); @@ -6913,7 +7026,8 @@ static void ggml_compute_forward_mul_mat_q_f32( ne11, ne01, ne10, 1.0f, y, ne10, x, ne00, - 0.0f, d, ne01); + 0.0f, d, ne01, + type); } } @@ -8875,6 +8989,111 @@ static void ggml_compute_forward_flash_ff( } } +// ggml_compute_forward_map_unary + +static void ggml_compute_forward_map_unary_f32( + const struct ggml_compute_params * params, + const struct ggml_tensor * src0, + struct ggml_tensor * dst, + const ggml_unary_op_f32_t fun) { + GGML_ASSERT(ggml_are_same_shape(src0, dst)); + + if (params->type == GGML_TASK_INIT || params->type == GGML_TASK_FINALIZE) { + return; + } + + const int n = ggml_nrows(src0); + const int nc = src0->ne[0]; + + assert( dst->nb[0] == sizeof(float)); + assert(src0->nb[0] == sizeof(float)); + + for (int i = 0; i < n; i++) { + fun(nc, + (float *) ((char *) dst->data + i*( dst->nb[1])), + (float *) ((char *) src0->data + i*(src0->nb[1]))); + } +} + + +static void ggml_compute_forward_map_unary( + const struct ggml_compute_params * params, + const struct ggml_tensor * src0, + struct ggml_tensor * dst, + const ggml_unary_op_f32_t fun) { + switch (src0->type) { + case GGML_TYPE_F32: + { + ggml_compute_forward_map_unary_f32(params, src0, dst, fun); + } break; + case GGML_TYPE_Q4_0: + case GGML_TYPE_Q4_1: + case GGML_TYPE_I8: + case GGML_TYPE_I16: + case GGML_TYPE_I32: + case GGML_TYPE_F16: + case GGML_TYPE_COUNT: + { + GGML_ASSERT(false); + } break; + } +} + +// ggml_compute_forward_map_binary + +static void ggml_compute_forward_map_binary_f32( + const struct ggml_compute_params * params, + const struct ggml_tensor * src0, + const struct ggml_tensor * src1, + struct ggml_tensor * dst, + const ggml_binary_op_f32_t fun) { + assert(params->ith == 0); + assert(ggml_are_same_shape(src0, src1) && ggml_are_same_shape(src0, dst)); + + if (params->type == GGML_TASK_INIT || params->type == GGML_TASK_FINALIZE) { + return; + } + + const int n = ggml_nrows(src0); + const int nc = src0->ne[0]; + + assert( dst->nb[0] == sizeof(float)); + assert(src0->nb[0] == sizeof(float)); + assert(src1->nb[0] == sizeof(float)); + + for (int i = 0; i < n; i++) { + fun(nc, + (float *) ((char *) dst->data + i*( dst->nb[1])), + (float *) ((char *) src0->data + i*(src0->nb[1])), + (float *) ((char *) src1->data + i*(src1->nb[1]))); + } +} + + +static void ggml_compute_forward_map_binary( + const struct ggml_compute_params * params, + const struct ggml_tensor * src0, + const struct ggml_tensor * src1, + struct ggml_tensor * dst, + const ggml_binary_op_f32_t fun) { + switch (src0->type) { + case GGML_TYPE_F32: + { + ggml_compute_forward_map_binary_f32(params, src0, src1, dst, fun); + } break; + case GGML_TYPE_Q4_0: + case GGML_TYPE_Q4_1: + case GGML_TYPE_I8: + case GGML_TYPE_I16: + case GGML_TYPE_I32: + case GGML_TYPE_F16: + case GGML_TYPE_COUNT: + { + GGML_ASSERT(false); + } break; + } +} + ///////////////////////////////// static void ggml_compute_forward(struct ggml_compute_params * params, struct ggml_tensor * tensor) { @@ -9024,6 +9243,18 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_flash_ff(params, tensor->src0, tensor->src1, tensor->opt[0], tensor->opt[1], tensor->opt[2], tensor); } break; + case GGML_OP_MAP_UNARY: + { + const ggml_unary_op_f32_t fun = *((ggml_unary_op_f32_t *)tensor->opt[0]->data); + ggml_compute_forward_map_unary(params, tensor->src0, tensor, fun); + } + break; + case GGML_OP_MAP_BINARY: + { + const ggml_binary_op_f32_t fun = *((ggml_binary_op_f32_t *)tensor->opt[0]->data); + ggml_compute_forward_map_binary(params, tensor->src0, tensor->src1, tensor, fun); + } + break; case GGML_OP_NONE: { // nop @@ -9283,6 +9514,11 @@ static void ggml_compute_backward(struct ggml_context * ctx, struct ggml_tensor { GGML_ASSERT(false); // not supported } break; + case GGML_OP_MAP_UNARY: + case GGML_OP_MAP_BINARY: + { + GGML_ASSERT(false); // not supported + } break; case GGML_OP_NONE: { // nop @@ -9775,6 +10011,11 @@ void ggml_graph_compute(struct ggml_context * ctx, struct ggml_cgraph * cgraph) work_size = MAX(work_size, cur); } break; + case GGML_OP_MAP_UNARY: + case GGML_OP_MAP_BINARY: + { + node->n_tasks = 1; + } break; case GGML_OP_NONE: { node->n_tasks = 1; diff --git a/ggml.h b/ggml.h index c06c09e060db..617298a95536 100644 --- a/ggml.h +++ b/ggml.h @@ -253,6 +253,9 @@ enum ggml_op { GGML_OP_FLASH_ATTN, GGML_OP_FLASH_FF, + GGML_OP_MAP_UNARY, + GGML_OP_MAP_BINARY, + GGML_OP_COUNT, }; @@ -351,6 +354,8 @@ int ggml_blck_size (enum ggml_type type); size_t ggml_type_size (enum ggml_type type); // size in bytes for all elements in a block float ggml_type_sizef(enum ggml_type type); // ggml_type_size()/ggml_blck_size() as float +const char * ggml_type_name(enum ggml_type type); + size_t ggml_element_size(const struct ggml_tensor * tensor); struct ggml_context * ggml_init(struct ggml_init_params params); @@ -652,6 +657,21 @@ struct ggml_tensor * ggml_flash_ff( struct ggml_tensor * c0, struct ggml_tensor * c1); +// Mapping operations +typedef void (*ggml_unary_op_f32_t)(const int, float *, const float *); +typedef void (*ggml_binary_op_f32_t)(const int, float *, const float *, const float *); + +struct ggml_tensor * ggml_map_unary_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + const ggml_unary_op_f32_t fun); + +struct ggml_tensor * ggml_map_binary_f32( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + const ggml_binary_op_f32_t fun); + // // automatic differentiation // diff --git a/ggml_blas_adapter.c b/ggml_blas_adapter.c index d41b001c9c21..a276faf58c9a 100644 --- a/ggml_blas_adapter.c +++ b/ggml_blas_adapter.c @@ -18,15 +18,65 @@ cl_platform_id platform; cl_device_id device; cl_context context; cl_command_queue queue; +cl_program program; +cl_kernel kernel_q4_0, kernel_q4_1; bool cl_initialized = false; -static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE trans_a, const enum CBLAS_TRANSPOSE trans_b, const int m, const int n, const int k, const float alpha, const float *host_a, const int lda, const float *host_b, const int ldb, const float beta, float *host_c, const int ldc) { +// Function taken from https://github.com/rsnemmen/OpenCL-examples/blob/master/add_numbers/add_numbers.c +cl_program build_program(cl_context ctx, cl_device_id dev, const char* filename) { + + cl_program program; + FILE *program_handle; + char *program_buffer, *program_log; + size_t program_size, log_size; + int err; + + program_handle = fopen(filename, "r"); + if(program_handle == NULL) { + perror("OpenCL kernel file not found"); + exit(1); + } + fseek(program_handle, 0, SEEK_END); + program_size = ftell(program_handle); + rewind(program_handle); + program_buffer = (char*)malloc(program_size + 1); + program_buffer[program_size] = '\0'; + fread(program_buffer, sizeof(char), program_size, program_handle); + fclose(program_handle); + + program = clCreateProgramWithSource(ctx, 1, + (const char**)&program_buffer, &program_size, &err); + if(err < 0) { + perror("OpenCL error creating program"); + exit(1); + } + free(program_buffer); + + err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + if(err < 0) { + + clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, + 0, NULL, &log_size); + program_log = (char*) malloc(log_size + 1); + program_log[log_size] = '\0'; + clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, + log_size + 1, program_log, NULL); + printf("%s\n", program_log); + free(program_log); + exit(1); + } + + return program; +} + +static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE trans_a, const enum CBLAS_TRANSPOSE trans_b, const int m, const int n, const int k, const float alpha, const void *host_a, const int lda, const float *host_b, const int ldb, const float beta, float *host_c, const int ldc, const int btype) { cl_int err = 0; - cl_event events[3]; + cl_event events[4]; events[0] = NULL; events[1] = NULL; events[2] = NULL; + events[3] = NULL; if (!cl_initialized) { char * KCPP_CLBLAST_PLATFORM = getenv("KCPP_CLBLAST_PLATFORM"); @@ -65,33 +115,82 @@ static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS free(platforms); free(devices); + program = build_program(context, device, "ggml_clblast_dequant.cl"); + + // Prepare dequantize kernels + kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err); + if(err < 0) { + printf("Error creating OpenCL dequantize q4_0 kernel: %d\n", err); + fflush(stdout); + }; + kernel_q4_1 = clCreateKernel(program, "dequantize_row_q4_1", &err); + if(err < 0) { + printf("Error creating OpenCL dequantize q4_1 kernel: %d\n", err); + fflush(stdout); + }; + cl_initialized = true; } + bool dequant = (btype == 2 || btype == 3); + cl_kernel kernel = btype == 2 ? kernel_q4_0 : kernel_q4_1; + + size_t global = n * k, local = 16, qb_size; + cl_mem cl_buffer_a, cl_buffer_qb, cl_buffer_b, cl_buffer_c; + // Prepare buffers - cl_mem cl_buffer_a = clCreateBuffer(context, CL_MEM_READ_ONLY, m*k*sizeof(float), NULL, &err); + cl_buffer_a = clCreateBuffer(context, CL_MEM_READ_ONLY, m*k*sizeof(float), NULL, &err); if (err != CL_SUCCESS) { printf("Error creating OpenCL Buffer A: %d\n", err); fflush(stdout); } - cl_mem cl_buffer_b = clCreateBuffer(context, CL_MEM_READ_ONLY, n*k*sizeof(float), NULL, &err); + if (dequant) { + qb_size = global * (sizeof(float) * (btype == 2 ? 1 : 2) + 16) / 32; + cl_buffer_qb = clCreateBuffer(context, CL_MEM_READ_ONLY, qb_size, NULL, &err); + if (err != CL_SUCCESS) { + printf("Error creating OpenCL Buffer QB: %d\n", err); + fflush(stdout); + } + } + cl_buffer_b = clCreateBuffer(context, CL_MEM_READ_WRITE, n*k*sizeof(float), NULL, &err); if (err != CL_SUCCESS) { printf("Error creating OpenCL Buffer B: %d\n", err); fflush(stdout); } - cl_mem cl_buffer_c = clCreateBuffer(context, CL_MEM_READ_WRITE, m*n*sizeof(float), NULL, &err); + cl_buffer_c = clCreateBuffer(context, CL_MEM_READ_WRITE, m*n*sizeof(float), NULL, &err); if (err != CL_SUCCESS) { printf("Error creating OpenCL Buffer C: %d\n", err); fflush(stdout); } + if (dequant) { + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_buffer_qb); + err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &cl_buffer_b); + if(err < 0) { + printf("Error setting OpenCL kernel args: %d\n", err); + fflush(stdout); + } + clEnqueueWriteBuffer(queue, cl_buffer_qb, CL_FALSE, 0, qb_size, host_b, 0, NULL, events + 1); + } else { + clEnqueueWriteBuffer(queue, cl_buffer_b, CL_FALSE, 0, n*k*sizeof(float), host_b, 0, NULL, events + 1); + } + clEnqueueWriteBuffer(queue, cl_buffer_a, CL_FALSE, 0, m*k*sizeof(float), host_a, 0, NULL, events); - clEnqueueWriteBuffer(queue, cl_buffer_b, CL_FALSE, 0, n*k*sizeof(float), host_b, 0, NULL, events + 1); clEnqueueWriteBuffer(queue, cl_buffer_c, CL_FALSE, 0, m*n*sizeof(float), host_c, 0, NULL, events + 2); - clWaitForEvents(3, events); + if (dequant) { + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 1, events + 1, events + 3); + if(err < 0) { + printf("Error enqueueing OpenCL dequantize kernel: %d\n", err); + fflush(stdout); + } + } + clWaitForEvents(dequant ? 4 : 3, events); clReleaseEvent(events[0]); clReleaseEvent(events[1]); clReleaseEvent(events[2]); + if (dequant) { + clReleaseEvent(events[3]); + } // Call the SGEMM routine. CLBlastStatusCode status = CLBlastSgemm(order, @@ -114,20 +213,22 @@ static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS } clReleaseMemObject(cl_buffer_a); + if (dequant) { + clReleaseMemObject(cl_buffer_qb); + } clReleaseMemObject(cl_buffer_b); clReleaseMemObject(cl_buffer_c); } - #endif #endif #if defined(GGML_USE_ACCELERATE) || defined(GGML_USE_OPENBLAS) #if GGML_USE_CLBLAST -#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc) ({\ -ggml_cl_sgemm_wrapper(Order, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, ldc);\ +#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc, btype) ({\ +ggml_cl_sgemm_wrapper(Order, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, ldc, btype);\ }) #else -#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc) ({\ +#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc, btype) ({\ cblas_sgemm(Order, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, ldc);\ }) #endif diff --git a/ggml_clblast_dequant.cl b/ggml_clblast_dequant.cl new file mode 100644 index 000000000000..87bdfe9bf25e --- /dev/null +++ b/ggml_clblast_dequant.cl @@ -0,0 +1,41 @@ +struct __attribute__ ((packed)) block_q4_0 +{ + float d; + uchar qs[16]; +}; + +__kernel void dequantize_row_q4_0(__global struct block_q4_0* blocks, __global float* result) { + uint i, l; + i = get_global_id(0) / 32; + l = get_local_id(0); + + float d = blocks[i].d; + + uchar vi = blocks[i].qs[l]; + + uint index = i*32 + l*2; + result[index + 0] = ((vi & 0xf) - 8)*d; + result[index + 1] = ((vi >> 4) - 8)*d; +} + +struct __attribute__ ((packed)) block_q4_1 +{ + float d; + float m; + uchar qs[16]; +}; + +__kernel void dequantize_row_q4_1(__global struct block_q4_1* blocks, __global float* result) { + uint i, l; + i = get_global_id(0) / 32; + l = get_local_id(0); + + float d = blocks[i].d; + float m = blocks[i].m; + + uchar vi = blocks[i].qs[l]; + + uint index = i*32 + l*2; + result[index + 0] = (vi & 0xf) * d + m; + result[index + 1] = (vi >> 4) * d + m; +} diff --git a/gpttype_adapter.cpp b/gpttype_adapter.cpp index 544c47bb8332..9c2a6d19facb 100644 --- a/gpttype_adapter.cpp +++ b/gpttype_adapter.cpp @@ -30,6 +30,7 @@ static int n_past = 0; static int n_threads = 4; static int n_batch = 8; static bool useSmartContext = false; +static int blasbatchsize = 512; static std::string modelname; static std::vector last_n_tokens; static std::vector current_context_tokens; @@ -53,6 +54,7 @@ ModelLoadResult gpttype_load_model(const load_model_inputs inputs, FileFormat in n_batch = params.n_batch = inputs.batch_size; modelname = params.model = inputs.model_filename; useSmartContext = inputs.use_smartcontext; + blasbatchsize = inputs.blasbatchsize; params.memory_f16 = inputs.f16_kv; params.n_ctx = inputs.max_context_length; model_v1.hparams.n_ctx = model_v2.hparams.n_ctx = model_gpt2_v1.hparams.n_ctx = model_gpt2_v2.hparams.n_ctx = params.n_ctx; @@ -208,7 +210,7 @@ generation_outputs gpttype_generate(const generation_inputs inputs, generation_o int original_threads = params.n_threads; if (blasmode) { - params.n_batch = 512; //received reports of 1024 and above crashing on some models + params.n_batch = blasbatchsize; //received reports of 1024 and above crashing on some models params.n_threads = 1; } diff --git a/koboldcpp.py b/koboldcpp.py index 86978534d228..6ac9911b98d3 100644 --- a/koboldcpp.py +++ b/koboldcpp.py @@ -17,7 +17,8 @@ class load_model_inputs(ctypes.Structure): ("n_parts_overwrite", ctypes.c_int), ("use_mmap", ctypes.c_bool), ("use_smartcontext", ctypes.c_bool), - ("clblast_info", ctypes.c_int)] + ("clblast_info", ctypes.c_int), + ("blasbatchsize", ctypes.c_int)] class generation_inputs(ctypes.Structure): _fields_ = [("seed", ctypes.c_int), @@ -66,7 +67,7 @@ def init_library(): handle.generate.argtypes = [generation_inputs, ctypes.c_wchar_p] #apparently needed for osx to work. i duno why they need to interpret it that way but whatever handle.generate.restype = generation_outputs -def load_model(model_filename,batch_size=8,max_context_length=512,n_parts_overwrite=-1,threads=6,use_mmap=False,use_smartcontext=False): +def load_model(model_filename,batch_size=8,max_context_length=512,n_parts_overwrite=-1,threads=6,use_mmap=False,use_smartcontext=False,blasbatchsize=512): inputs = load_model_inputs() inputs.model_filename = model_filename.encode("UTF-8") inputs.batch_size = batch_size @@ -76,6 +77,7 @@ def load_model(model_filename,batch_size=8,max_context_length=512,n_parts_overwr inputs.f16_kv = True inputs.use_mmap = use_mmap inputs.use_smartcontext = use_smartcontext + inputs.blasbatchsize = blasbatchsize clblastids = 0 if args.useclblast: clblastids = 100 + int(args.useclblast[0])*10 + int(args.useclblast[1]) @@ -355,8 +357,10 @@ def main(args): print("Overriding thread count, using " + str(args.threads) + " threads instead.") init_library() # Note: if blas does not exist and is enabled, program will crash. - ggml_selected_file = args.model_file embedded_kailite = None + ggml_selected_file = args.model_param + if not ggml_selected_file: + ggml_selected_file = args.model if not ggml_selected_file: #give them a chance to pick a file print("For command line arguments, please refer to --help") @@ -386,7 +390,7 @@ def main(args): mdl_nparts = sum(1 for n in range(1, 9) if os.path.exists(f"{ggml_selected_file}.{n}")) + 1 modelname = os.path.abspath(ggml_selected_file) print(f"Loading model: {modelname} \n[Parts: {mdl_nparts}, Threads: {args.threads}, SmartContext: {args.smartcontext}]") - loadok = load_model(modelname,8,maxctx,mdl_nparts,args.threads,(not args.nommap),args.smartcontext) + loadok = load_model(modelname,8,maxctx,mdl_nparts,args.threads,(not args.nommap),args.smartcontext,args.blasbatchsize) print("Load Model OK: " + str(loadok)) if not loadok: @@ -401,8 +405,8 @@ def main(args): except: print("Could not find Kobold Lite. Embedded Kobold Lite will not be available.") - if args.l_port!=defaultport: - args.port = args.l_port + if args.port_param!=defaultport: + args.port = args.port_param print(f"Starting Kobold HTTP Server on port {args.port}") epurl = "" if args.host=="": @@ -415,12 +419,14 @@ def main(args): RunServerMultiThreaded(args.host, args.port, embedded_kailite) if __name__ == '__main__': - print("Welcome to KoboldCpp - Version 1.7 (Wordstopper fork)") # just update version manually + print("Welcome to KoboldCpp - Version 1.7 (Wordstopper fork)") parser = argparse.ArgumentParser(description='Kobold llama.cpp server') - parser.add_argument("model_file", help="Model file to load", nargs="?") + modelgroup = parser.add_mutually_exclusive_group() #we want to be backwards compatible with the unnamed positional args + modelgroup.add_argument("--model", help="Model file to load", nargs="?") + modelgroup.add_argument("model_param", help="Model file to load (positional)", nargs="?") portgroup = parser.add_mutually_exclusive_group() #we want to be backwards compatible with the unnamed positional args portgroup.add_argument("--port", help="Port to listen on", default=defaultport, type=int, action='store') - portgroup.add_argument("l_port", help="Port to listen on (deprecated)", default=defaultport, nargs="?", type=int, action='store') + portgroup.add_argument("port_param", help="Port to listen on (positional)", default=defaultport, nargs="?", type=int, action='store') parser.add_argument("--host", help="Host IP to listen on. If empty, all routable interfaces are accepted.", default="") #os.environ["OMP_NUM_THREADS"] = '12' @@ -431,6 +437,7 @@ def main(args): default_threads = (physical_core_limit if physical_core_limit<=3 else max(3,physical_core_limit-1)) parser.add_argument("--threads", help="Use a custom number of threads if specified. Otherwise, uses an amount based on CPU cores", type=int, default=default_threads) parser.add_argument("--psutil_set_threads", help="Experimental flag. If set, uses psutils to determine thread count based on physical cores.", action='store_true') + parser.add_argument("--blasbatchsize", help="Sets the batch size used in BLAS processing (default 512)", type=int,choices=[128,256,512,1024], default=512) parser.add_argument("--stream", help="Uses pseudo streaming", action='store_true') parser.add_argument("--smartcontext", help="Reserving a portion of context to try processing less frequently.", action='store_true') parser.add_argument("--nommap", help="If set, do not use mmap to load newer models", action='store_true') diff --git a/llama.cpp b/llama.cpp index 1eba38d4c1a9..2c931480ab44 100644 --- a/llama.cpp +++ b/llama.cpp @@ -269,16 +269,6 @@ static std::string llama_format_tensor_shape(const std::vector & ne) { return ret; } -static const char * llama_format_type(enum ggml_type type) { - switch (type) { - case GGML_TYPE_F32: return "f32"; - case GGML_TYPE_F16: return "f16"; - case GGML_TYPE_Q4_0: return "q4_0"; - case GGML_TYPE_Q4_1: return "q4_1"; - default: LLAMA_ASSERT(false); - } -} - static size_t llama_calc_tensor_size(const std::vector & ne, enum ggml_type type) { size_t size = ggml_type_size(type); for (uint32_t dim : ne) { @@ -1589,7 +1579,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s printf("[%zu/%zu] %36s - %s, type = %6s, ", ++idx, model_loader->tensors_map.tensors.size(), tensor.name.c_str(), llama_format_tensor_shape(tensor.ne).c_str(), - llama_format_type(tensor.type)); + ggml_type_name(tensor.type)); // This used to be a regex, but has an extreme cost to compile times. bool quantize = tensor.name.rfind("weight") == tensor.name.size() - 6; // ends with 'weight'? @@ -1622,7 +1612,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s f32_data[i] = ggml_fp16_to_fp32(f16_data[i]); } } else { - throw format("type %s unsupported for integer quantization", llama_format_type(tensor.type)); + throw format("type %s unsupported for integer quantization", ggml_type_name(tensor.type)); } printf("quantizing .. "); diff --git a/llama_adapter.cpp b/llama_adapter.cpp index b9d64b969e05..900cf27ed960 100644 --- a/llama_adapter.cpp +++ b/llama_adapter.cpp @@ -28,6 +28,7 @@ static int n_past = 0; static int n_threads = 4; static int n_batch = 8; static bool useSmartContext = false; +static int blasbatchsize = 512; static std::string modelname; static llama_context *ctx; static std::vector last_n_tokens; @@ -44,6 +45,7 @@ bool llama_load_model(const load_model_inputs inputs, FileFormat in_file_format) n_batch = inputs.batch_size; modelname = inputs.model_filename; useSmartContext = inputs.use_smartcontext; + blasbatchsize = inputs.blasbatchsize; ctx_params.n_ctx = inputs.max_context_length; ctx_params.n_parts = -1;//inputs.n_parts_overwrite; @@ -143,7 +145,7 @@ generation_outputs llama_generate(const generation_inputs inputs, generation_out int original_threads = params.n_threads; if (blasmode) { - params.n_batch = 512; //received reports of 1024 and above crashing on some models + params.n_batch = blasbatchsize; //received reports of 1024 and above crashing on some models params.n_threads = 1; } @@ -160,7 +162,7 @@ generation_outputs llama_generate(const generation_inputs inputs, generation_out double time1 = 0, time2 = 0; unsigned int embd_inp_size = embd_inp.size(); printf("\n"); - + std::ifstream config_file("wordstoppers.txt"); // Open config file for reading std::vector search_strings; @@ -178,7 +180,7 @@ generation_outputs llama_generate(const generation_inputs inputs, generation_out std::cerr << "Error: Failed to find wordstoppers.txt file." << std::endl; } - + while (remaining_tokens > 0) { llama_token id = 0; diff --git a/requirements.txt b/requirements.txt index f3944951a486..6c32cbd047b8 100644 --- a/requirements.txt +++ b/requirements.txt @@ -1,2 +1,2 @@ numpy==1.24 -sentencepiece==0.1.97 +sentencepiece==0.1.98