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);