backport MM256_SET_M128I to ggml_v2, updated lite, added support for …

…selecting the GPU for cublas

expose.cpp

@@ -47,14 +47,14 @@ extern "C"
         }
 
         //first digit is whether configured, second is platform, third is devices
-        int parseinfo = inputs.clblast_info;
+        int cl_parseinfo = inputs.clblast_info;
 
-        std::string usingclblast = "GGML_OPENCL_CONFIGURED="+std::to_string(parseinfo>0?1:0);
+        std::string usingclblast = "GGML_OPENCL_CONFIGURED="+std::to_string(cl_parseinfo>0?1:0);
         putenv((char*)usingclblast.c_str());
 
-        parseinfo = parseinfo%100; //keep last 2 digits
-        int platform = parseinfo/10;
-        int devices = parseinfo%10;
+        cl_parseinfo = cl_parseinfo%100; //keep last 2 digits
+        int platform = cl_parseinfo/10;
+        int devices = cl_parseinfo%10;
         platformenv = "GGML_OPENCL_PLATFORM="+std::to_string(platform);
         deviceenv = "GGML_OPENCL_DEVICE="+std::to_string(devices);
         putenv((char*)platformenv.c_str());

expose.h

@@ -30,6 +30,7 @@ struct load_model_inputs
     const bool use_smartcontext;
     const bool unban_tokens;
     const int clblast_info = 0;
+    const int cublas_info = 0;
     const int blasbatchsize = 512;
     const int debugmode = 0;
     const int forceversion = 0;

ggml.c

@@ -588,8 +588,9 @@ static const size_t CACHE_LINE_SIZE_F32 = CACHE_LINE_SIZE/sizeof(float);
 //
 // quantization
 //
-
+#ifndef MM256_SET_M128I
 #define MM256_SET_M128I(a, b) _mm256_insertf128_si256(_mm256_castsi128_si256(b), (a), 1)
+#endif
 
 #if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) || defined(__SSSE3__)
 // multiply int8_t, add results pairwise twice

gpttype_adapter.cpp

@@ -347,7 +347,16 @@ ModelLoadResult gpttype_load_model(const load_model_inputs inputs, FileFormat in
     //this is used for the mem_per_token eval, openblas needs more RAM
     bool use_scratch = ggml_cpu_has_gpublas();
 
+    int cu_parseinfo_maindevice = inputs.cublas_info<0?0:inputs.cublas_info;
+
     printf("System Info: %s\n", llama_print_system_info());
+    #if defined(GGML_USE_CUBLAS)
+    if(ggml_cpu_has_gpublas() && cu_parseinfo_maindevice>0)
+    {
+        printf("CUBLAS: Set main device to %d\n",cu_parseinfo_maindevice);
+        ggml_cuda_set_main_device(cu_parseinfo_maindevice);
+    }
+    #endif
     SetQuantsUnshuffled(false);
     if(file_format == FileFormat::GGML || file_format == FileFormat::GGHF || file_format == FileFormat::GGJT || file_format == FileFormat::GGJT_2)
     {
@@ -412,6 +421,7 @@ ModelLoadResult gpttype_load_model(const load_model_inputs inputs, FileFormat in
         llama_ctx_params.use_mmap = inputs.use_mmap;
         llama_ctx_params.use_mlock = inputs.use_mlock;
         llama_ctx_params.n_gpu_layers = inputs.gpulayers;
+        llama_ctx_params.main_gpu = cu_parseinfo_maindevice;
 
         llama_ctx_v3 = llama_init_from_file(modelname.c_str(), llama_ctx_params);
 

koboldcpp.py

@@ -30,6 +30,7 @@ class load_model_inputs(ctypes.Structure):
                 ("use_smartcontext", ctypes.c_bool),
                 ("unban_tokens", ctypes.c_bool),
                 ("clblast_info", ctypes.c_int),
+                ("cublas_info", ctypes.c_int),
                 ("blasbatchsize", ctypes.c_int),
                 ("debugmode", ctypes.c_int),
                 ("forceversion", ctypes.c_int),
@@ -111,7 +112,7 @@ def init_library():
         else:
             print("Attempting to use CLBlast library for faster prompt ingestion. A compatible clblast will be required.")
             use_clblast = True
-    elif (args.usecublas and args.usecublas!=""):
+    elif (args.usecublas is not None):
         if not file_exists(lib_cublas):
             print("Warning: CuBLAS library file not found. Non-BLAS library will be used.")
         else:
@@ -166,7 +167,7 @@ def load_model(model_filename):
     inputs.batch_size = 8
     inputs.max_context_length = maxctx #initial value to use for ctx, can be overwritten
     inputs.threads = args.threads
-    inputs.low_vram = (True if args.usecublas=="lowvram" else False)
+    inputs.low_vram = (True if (args.usecublas and "lowvram" in args.usecublas) else False)
     inputs.blasthreads = args.blasthreads
     inputs.f16_kv = True
     inputs.use_mmap = (not args.nommap)
@@ -187,6 +188,11 @@ def load_model(model_filename):
     if args.useclblast:
         clblastids = 100 + int(args.useclblast[0])*10 + int(args.useclblast[1])
     inputs.clblast_info = clblastids
+    inputs.cublas_info = 0
+    if (args.usecublas and "1" in args.usecublas):
+        inputs.cublas_info = 1
+    elif (args.usecublas and "2" in args.usecublas):
+        inputs.cublas_info = 2
     inputs.executable_path = (getdirpath()+"/").encode("UTF-8")
     inputs.debugmode = args.debugmode
     ret = handle.load_model(inputs)
@@ -805,7 +811,7 @@ def changerunmode(a,b,c):
     makeslider(quick_tab, "Context Size:", contextsize_text, context_var, 0, len(contextsize_text)-1, 30, set=2)
 
     # load model
-    makefileentry(quick_tab, "Model:", "Select Model File", model_var, 40, 170)
+    makefileentry(quick_tab, "Model:", "Select GGML Model File", model_var, 40, 170)
 
     # Hardware Tab
     hardware_tab = tabcontent["Hardware"]
@@ -867,7 +873,7 @@ def togglemiro(a,b,c):
     # Model Tab
     model_tab = tabcontent["Model"]
 
-    makefileentry(model_tab, "Model:", "Select Model File", model_var, 1)
+    makefileentry(model_tab, "Model:", "Select GGML Model File", model_var, 1)
     makefileentry(model_tab, "Lora:", "Select Lora File",lora_var, 3)
     makefileentry(model_tab, "Lora Base:", "Select Lora Base File", lora_base_var, 5)
 
@@ -947,7 +953,7 @@ def switch_old_gui():
         if runopts_var.get() == runopts[1]:
             args.useclblast = [[0,0], [1,0], [0,1]][int(gpu_choice_var.get())-1]
         if runopts_var.get() == runopts[2]:
-            args.usecublas = "lowvram" if lowvram_var.get() == 1 else "normal"
+            args.usecublas = ["lowvram"] if lowvram_var.get() == 1 else ["normal"]
         if gpulayers_var.get():
             args.gpulayers = int(gpulayers_var.get())
         if runopts_var.get()==runopts[3]:
@@ -1094,7 +1100,7 @@ def onDropdownChange(event):
         if selrunchoice==runopts[3]:
             args.useclblast = [0,1]
         if selrunchoice==runopts[4]:
-            args.usecublas = True
+            args.usecublas = ["normal"]
         if selrunchoice==runopts[5]:
             args.noblas = True
         if selrunchoice==runopts[6]:
@@ -1290,7 +1296,8 @@ def main(args):
     parser.add_argument("--blasbatchsize", help="Sets the batch size used in BLAS processing (default 512). Setting it to -1 disables BLAS mode, but keeps other benefits like GPU offload.", type=int,choices=[-1,32,64,128,256,512,1024], default=512)
     parser.add_argument("--stream", help="Uses streaming when generating tokens. Only for the Kobold Lite UI.", action='store_true')
     parser.add_argument("--smartcontext", help="Reserving a portion of context to try processing less frequently.", action='store_true')
-    parser.add_argument("--unbantokens", help="Normally, KoboldAI prevents certain tokens such as EOS and Square Brackets. This flag unbans them.", action='store_true')
+    parser.add_argument("--unbantokens", help="Normally, KoboldAI prevents the EOS token from being generated. This flag unbans it.", action='store_true')
+    parser.add_argument("--bantokens", help="You can manually specify a list of token IDs that the AI cannot use.", metavar=('[elements]'), nargs='+')
     parser.add_argument("--usemirostat", help="Experimental! Replaces your samplers with mirostat. Takes 3 params = [type(0/1/2), tau(5.0), eta(0.1)].",metavar=('[type]', '[tau]', '[eta]'), type=float, nargs=3)
     parser.add_argument("--forceversion", help="If the model file format detection fails (e.g. rogue modified model) you can set this to override the detected format (enter desired version, e.g. 401 for GPTNeoX-Type2).",metavar=('[version]'), type=int, default=0)
     parser.add_argument("--nommap", help="If set, do not use mmap to load newer models", action='store_true')
@@ -1302,7 +1309,7 @@ def main(args):
     compatgroup = parser.add_mutually_exclusive_group()
     compatgroup.add_argument("--noblas", help="Do not use OpenBLAS for accelerated prompt ingestion", action='store_true')
     compatgroup.add_argument("--useclblast", help="Use CLBlast for GPU Acceleration. Must specify exactly 2 arguments, platform ID and device ID (e.g. --useclblast 1 0).", type=int, choices=range(0,9), nargs=2)
-    compatgroup.add_argument("--usecublas", help="Use CuBLAS for GPU Acceleration. Requires Nvidia GPU. Select lowvram to not allocate VRAM scratch buffer.", default='', const='normal', nargs='?', choices=['normal', 'lowvram'])
+    compatgroup.add_argument("--usecublas", help="Use CuBLAS for GPU Acceleration. Requires Nvidia GPU. Select lowvram to not allocate VRAM scratch buffer. Enter a number after to select a different main GPU.", nargs='*',metavar=('[lowvram|normal] [main GPU ID]'), choices=['normal', 'lowvram', '0', '1', '2'])
     parser.add_argument("--gpulayers", help="Set number of layers to offload to GPU when using GPU. Requires GPU.",metavar=('[GPU layers]'), type=int, default=0)
     args = parser.parse_args()
     main(args)

otherarch/ggml_v2.c

@@ -472,6 +472,9 @@ static const size_t CACHE_LINE_SIZE_F32 = CACHE_LINE_SIZE/sizeof(float);
 //
 // quantization
 //
+#ifndef MM256_SET_M128I
+#define MM256_SET_M128I(a, b) _mm256_insertf128_si256(_mm256_castsi128_si256(b), (a), 1)
+#endif
 
 #if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) || defined(__SSSE3__)
 // multiply int8_t, add results pairwise twice
@@ -532,7 +535,7 @@ static inline __m256i bytes_from_bits_32(const uint8_t * x) {
 static inline __m256i bytes_from_nibbles_32(const uint8_t * rsi)
 {
     const __m128i tmp = _mm_loadu_si128((const __m128i *)rsi);
-    const __m256i bytes = _mm256_set_m128i(_mm_srli_epi16(tmp, 4), tmp);
+    const __m256i bytes = MM256_SET_M128I(_mm_srli_epi16(tmp, 4), tmp);
     const __m256i lowMask = _mm256_set1_epi8( 0xF );
     return _mm256_and_si256(lowMask, bytes);
 }
@@ -605,7 +608,7 @@ static inline __m256i bytes_from_bits_32(const uint8_t * x) {
     bytesh = _mm_or_si128(bytesh, bit_mask);
     bytesl = _mm_cmpeq_epi8(bytesl, _mm_set1_epi64x(-1));
     bytesh = _mm_cmpeq_epi8(bytesh, _mm_set1_epi64x(-1));
-    return _mm256_set_m128i(bytesh, bytesl);
+    return MM256_SET_M128I(bytesh, bytesl);
 }
 
 // Unpack 32 4-bit fields into 32 bytes
@@ -618,15 +621,15 @@ static inline __m256i bytes_from_nibbles_32(const uint8_t * rsi)
     const __m128i lowMask = _mm_set1_epi8(0xF);
     tmpl = _mm_and_si128(lowMask, tmpl);
     tmph = _mm_and_si128(lowMask, tmph);
-    return _mm256_set_m128i(tmph, tmpl);
+    return MM256_SET_M128I(tmph, tmpl);
 }
 
 // add int16_t pairwise and return as float vector
 static inline __m256 sum_i16_pairs_float(const __m128i xh, const __m128i xl) {
     const __m128i ones = _mm_set1_epi16(1);
     const __m128i summed_pairsl = _mm_madd_epi16(ones, xl);
     const __m128i summed_pairsh = _mm_madd_epi16(ones, xh);
-    const __m256i summed_pairs = _mm256_set_m128i(summed_pairsh, summed_pairsl);
+    const __m256i summed_pairs = MM256_SET_M128I(summed_pairsh, summed_pairsl);
     return _mm256_cvtepi32_ps(summed_pairs);
 }
 
@@ -2246,7 +2249,7 @@ static void ggml_v2_vec_dot_q4_0_q8_0(const int n, float * restrict s, const voi
         const __m128i i32_1 = mul_sum_i8_pairs(bx, by);
 
         // Convert int32_t to float
-        __m256 p = _mm256_cvtepi32_ps(_mm256_set_m128i(i32_0, i32_1));
+        __m256 p = _mm256_cvtepi32_ps(MM256_SET_M128I(i32_0, i32_1));
 
         // Apply the scale, and accumulate
         acc = _mm256_add_ps(_mm256_mul_ps( d, p ), acc);
@@ -2727,7 +2730,7 @@ static void ggml_v2_vec_dot_q5_0_q8_0(const int n, float * restrict s, const voi
         __m128i bxh = _mm256_extractf128_si256(bx, 1);
         bxl = _mm_or_si128(bxl, bxhil);
         bxh = _mm_or_si128(bxh, bxhih);
-        bx = _mm256_set_m128i(bxh, bxl);
+        bx = MM256_SET_M128I(bxh, bxl);
 
         const __m256i by = _mm256_loadu_si256((const __m256i *)y[i].qs);
 
@@ -2989,7 +2992,7 @@ static void ggml_v2_vec_dot_q5_1_q8_1(const int n, float * restrict s, const voi
         __m128i bxh = _mm256_extractf128_si256(bx, 1);
         bxl = _mm_or_si128(bxl, bxhil);
         bxh = _mm_or_si128(bxh, bxhih);
-        bx = _mm256_set_m128i(bxh, bxl);
+        bx = MM256_SET_M128I(bxh, bxl);
 
         const __m256 dy = _mm256_broadcast_ss(&y[i].d);
         const __m256i by = _mm256_loadu_si256((const __m256i *)y[i].qs);
@@ -17417,7 +17420,7 @@ static void ggml_v2_vec_dot_q4_0_q8_0_v2(const int n, float * restrict s, const
         }
 
         // Convert int32_t to float
-        __m256 p = _mm256_cvtepi32_ps( _mm256_set_m128i( i32[0], i32[1] ));
+        __m256 p = _mm256_cvtepi32_ps( MM256_SET_M128I( i32[0], i32[1] ));
         // Apply the scale, and accumulate
         acc = _mm256_add_ps(_mm256_mul_ps( d, p ), acc);
     }
@@ -17686,7 +17689,7 @@ static void ggml_v2_vec_dot_q4_2_q8_0_v2(const int n, float * restrict s, const
 
         __m128i bx0 = bytes_from_nibbles_16(x[2*i + 0].qs);
         __m128i bx1 = bytes_from_nibbles_16(x[2*i + 1].qs);
-        __m256i bx = _mm256_set_m128i(bx1, bx0);
+        __m256i bx = MM256_SET_M128I(bx1, bx0);
 
         // Now we have a vector with bytes in [ 0 .. 15 ] interval. Offset them into [ -8 .. +7 ] interval.
         const __m256i off = _mm256_set1_epi8(8);
@@ -17819,7 +17822,7 @@ static void ggml_v2_vec_dot_q4_3_q8_1_v2(const int n, float * restrict s, const
 
         const __m128i bx0 = bytes_from_nibbles_16(x[2*i + 0].qs);
         const __m128i bx1 = bytes_from_nibbles_16(x[2*i + 1].qs);
-        const __m256i bx = _mm256_set_m128i(bx1, bx0);
+        const __m256i bx = MM256_SET_M128I(bx1, bx0);
 
         const __m256 dy = _mm256_broadcast_ss(&y[i].d);
         const __m256i by = _mm256_loadu_si256((const __m256i *)y[i].qs);