-
Notifications
You must be signed in to change notification settings - Fork 0
Expand file tree
/
Copy pathvec.cpp
More file actions
5128 lines (4570 loc) · 188 KB
/
vec.cpp
File metadata and controls
5128 lines (4570 loc) · 188 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
781
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
935
936
937
938
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980
981
982
983
984
985
986
987
988
989
990
991
992
993
994
995
996
997
998
999
1000
/*
* Curated by @PsyChip
* root@psychip.net
* April 2026
*
* vec - dead simple GPU-resident vector database
*
* Usage: vec <name> <dim[:format]> [port]
* vec deploy [port]
* vec --deploy=name:dim[:fmt][,...] [port]
* format: f32 (default), f16
*
* Creates an in-memory CUDA vector store.
* Listens on:
* - TCP port (default 1920)
* - Windows: Named pipe \\.\pipe\vec_<name>
* - Linux: Unix socket /tmp/vec_<name>.sock
*
* VEC 2.0 binary frame protocol — see PROTOCOL-2.0.md for the byte-exact spec.
* request: F0 <2B ns_len> [ns] <CMD> <2B label_len> [label] <4B body_len> [body]
* response: <1B status> <4B body_len> [body] ; status 0=OK, 1=ERR
* CMD: 01=push 02=query 04=get 06=update 07=delete 08=label
* 09=undo 0A=save 0D=cluster 0E=distinct 0F=represent
* 10=info 11=qid 13=set_data 14=get_data 15=exists
*
* GPU top-K kernel for datasets above 100K entries.
* Labels: filename-scheme, ≤2048 bytes, validated on write, lenient on load.
* Data: opaque blobs ≤100KB, requires a label.
*
* File format:
* .tensors [4B dim][4B count][4B deleted][1B fmt][count B alive][vectors][4B CRC32]
* .meta [4B count][per slot: 4B len + label bytes]
* .data [4B count][count B alive mask][per present slot: 4B len + bytes][4B CRC32]
*
* Features:
* - Read-only mode if file is not writable
* - Disk space check before save
* - File integrity check/repair (--check, --repair)
*
* Ctrl+C saves before exit.
*
* Build (Windows):
* nvcc -O2 -c vec_kernel.cu -o vec_kernel.obj <gencode flags>
* nvcc -O2 vec_kernel.obj vec.cpp -o vec.exe -lws2_32 <gencode flags>
*
* Build (Linux):
* nvcc -O2 -c vec_kernel.cu -o vec_kernel.o <gencode flags>
* nvcc -O2 vec_kernel.o vec.cpp -o vec -lpthread <gencode flags>
*/
/* ===================================================================== */
/* Platform includes */
/* ===================================================================== */
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#include <winsock2.h>
#include <ws2tcpip.h>
#pragma comment(lib, "ws2_32.lib")
#else
#include <sys/socket.h>
#include <sys/un.h>
#include <sys/file.h>
#include <sys/stat.h>
#include <sys/statvfs.h>
#include <sys/wait.h>
#include <netinet/in.h>
#include <arpa/inet.h>
#include <unistd.h>
#include <pthread.h>
#include <signal.h>
#include <dirent.h>
#include <glob.h>
#include <fcntl.h>
#include <errno.h>
#endif
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <float.h>
#include <stdint.h>
#include <unordered_map>
#include "xxhash.h"
/* ===================================================================== */
/* Constants */
/* ===================================================================== */
#define DEFAULT_PORT 1920
#define DEFAULT_TOP_K 10
#define GPU_TOP_K 10
#define GPU_SORT_THRESHOLD 100000
#define INITIAL_CAP 4096
#define MAX_LINE (1 << 24)
#define PIPE_BUF_SIZE (1 << 20)
#define FMT_F32 0
#define FMT_F16 1
/* VEC 2.0 binary frame protocol — see PROTOCOL-2.0.md */
#define BIN_MAGIC 0xF0
#define PROTOCOL_VERSION 0x02
#define CMD_PUSH 0x01
#define CMD_QUERY 0x02 /* unified PULL+CPULL */
/* 0x03 removed — was CPULL */
#define CMD_GET 0x04 /* unified GET+MGET */
/* 0x05 removed — was MGET */
#define CMD_UPDATE 0x06
#define CMD_DELETE 0x07
#define CMD_LABEL 0x08
#define CMD_UNDO 0x09
#define CMD_SAVE 0x0A
/* 0x0B, 0x0C reserved (do not reuse) */
#define CMD_CLUSTER 0x0D
#define CMD_DISTINCT 0x0E
#define CMD_REPRESENT 0x0F
#define CMD_INFO 0x10
#define CMD_QID 0x11 /* unified PID+CPID */
/* 0x12 removed — was CPID */
#define CMD_SET_DATA 0x13
#define CMD_GET_DATA 0x14
#define CMD_EXISTS 0x15
/* response envelope */
#define RESP_OK 0x00
#define RESP_ERR 0x01
/* shape mask bits */
#define SHAPE_VECTOR 0x01
#define SHAPE_LABEL 0x02
#define SHAPE_DATA 0x04
/* GET mode */
#define GET_MODE_SINGLE 0x00
#define GET_MODE_BATCH 0x01
/* metric */
#define METRIC_L2 0x00
#define METRIC_COSINE 0x01
#ifdef _WIN32
#define strtok_r strtok_s
#endif
/* ===================================================================== */
/* Elapsed-time helper */
/* ===================================================================== */
/* monotonic timestamp in ms since some arbitrary epoch */
static long long now_ms() {
#ifdef _WIN32
LARGE_INTEGER c, f;
QueryPerformanceCounter(&c);
QueryPerformanceFrequency(&f);
return (long long)((c.QuadPart * 1000LL) / f.QuadPart);
#else
struct timespec ts;
clock_gettime(CLOCK_MONOTONIC, &ts);
return (long long)ts.tv_sec * 1000LL + ts.tv_nsec / 1000000LL;
#endif
}
/* format elapsed ms into a compact human string: "23ms", "4.1s", "52m", "2h13m" */
static void format_elapsed(long long ms, char *out, int outsz) {
if (ms < 1000) {
snprintf(out, outsz, "%lldms", ms);
} else if (ms < 60000) {
snprintf(out, outsz, "%.1fs", ms / 1000.0);
} else if (ms < 3600000) {
long long m = ms / 60000;
long long s = (ms % 60000) / 1000;
snprintf(out, outsz, "%lldm%llds", m, s);
} else {
long long h = ms / 3600000;
long long m = (ms % 3600000) / 60000;
snprintf(out, outsz, "%lldh%lldm", h, m);
}
}
/* ===================================================================== */
/* CRC32 */
/* ===================================================================== */
static unsigned int crc32_tab[256];
static int crc32_ready = 0;
static void crc32_init() {
for (unsigned int i = 0; i < 256; i++) {
unsigned int c = i;
for (int j = 0; j < 8; j++) c = (c >> 1) ^ ((c & 1) ? 0xEDB88320 : 0);
crc32_tab[i] = c;
}
crc32_ready = 1;
}
static unsigned int crc32_update(unsigned int crc, const void *buf, size_t len) {
if (!crc32_ready) crc32_init();
const unsigned char *p = (const unsigned char *)buf;
crc = ~crc;
for (size_t i = 0; i < len; i++) crc = crc32_tab[(crc ^ p[i]) & 0xFF] ^ (crc >> 8);
return ~crc;
}
static const char CRC_C[] = "BCDFGHJKLMNPRSTVZ"; /* 17 */
static const char CRC_V[] = "AEIOU"; /* 5 */
#define CRC_SYL (17 * 5) /* 85 per syllable, 85^4 = 52M unique */
static void crc32_word(unsigned int crc, char *out) {
unsigned int v = crc;
for (int i = 0; i < 4; i++) {
int s = v % CRC_SYL;
v /= CRC_SYL;
*out++ = CRC_C[s / 5];
*out++ = CRC_V[s % 5];
}
*out = '\0';
}
/* ===================================================================== */
/* External kernels (vec_kernel.cu) */
/* ===================================================================== */
extern "C" {
void launch_l2_f32(const float *db, const float *query, float *dists, int n, int dim);
void launch_cos_f32(const float *db, const float *query, float *dists, int n, int dim);
void launch_l2_f16(const void *db, const void *query, float *dists, int n, int dim);
void launch_cos_f16(const void *db, const void *query, float *dists, int n, int dim);
void launch_f32_to_f16(const float *src, void *dst, int count);
void launch_topk(const float *d_dists, const unsigned char *d_alive, int n,
float *out_dists, int *out_ids);
}
/* ===================================================================== */
/* Globals */
/* ===================================================================== */
static char g_name[256];
static char g_filepath[512];
static int g_dim = 0;
static int g_port = DEFAULT_PORT;
static int g_fmt = FMT_F32;
static int g_elem_size = 4;
static const char *fmt_name(int fmt) { return fmt == FMT_F16 ? "f16" : "f32"; }
static int g_count = 0;
static int g_capacity = 0;
static void *d_vectors = NULL;
static void *d_query = NULL;
static float *d_dists = NULL;
static float *d_topk_dists = NULL;
static int *d_topk_ids = NULL;
static unsigned char *d_alive = NULL;
static float *d_staging = NULL;
static int d_staging_n = 0;
static float *h_dists = NULL;
static int *h_ids = NULL;
static unsigned char *g_alive = NULL;
static int g_alive_cap = 0;
static int g_deleted = 0;
/* labels: variable-length strings indexed by slot */
static char **g_labels = NULL; /* array of pointers, NULL = no label */
static int g_labels_cap = 0;
/* data blobs: variable-length opaque payloads indexed by slot, ≤ MAX_DATA_BYTES each */
#define MAX_DATA_BYTES 102400
#define MAX_LABEL_BYTES 2048
static unsigned char **g_blobs = NULL; /* array of pointers, NULL = no blob */
static unsigned int *g_blob_lens = NULL; /* 0 if g_blobs[i] == NULL */
static int g_blobs_cap = 0;
/* hash index: xxh64 of stored vector bytes -> slot. multimap absorbs the rare
* collision; lookup walks all entries for a hash and byte-compares each.
* tombstoned slots are erased from the map on DELETE/UNDO so EXISTS skips them. */
static std::unordered_multimap<uint64_t, uint32_t> g_hash_index;
static int g_dirty = 0;
static int g_readonly = 0;
static volatile int g_running = 1;
static volatile time_t g_last_write = 0;
#define AUTOSAVE_IDLE_SECS 60
#ifdef _WIN32
static HANDLE g_mutex = NULL;
static CRITICAL_SECTION g_req_mutex;
#define REQ_LOCK() EnterCriticalSection(&g_req_mutex)
#define REQ_UNLOCK() LeaveCriticalSection(&g_req_mutex)
#else
static char g_sockpath[512];
static int g_lockfd = -1;
static pthread_mutex_t g_req_mutex = PTHREAD_MUTEX_INITIALIZER;
#define REQ_LOCK() pthread_mutex_lock(&g_req_mutex)
#define REQ_UNLOCK() pthread_mutex_unlock(&g_req_mutex)
#endif
/* ===================================================================== */
/* CUDA helpers */
/* ===================================================================== */
static void save_to_file(int already_locked); /* forward declaration for CUDA_CHECK */
#define CUDA_CHECK(x) do { \
cudaError_t err = (x); \
if (err != cudaSuccess) { \
fprintf(stderr, "FATAL: CUDA error %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
save_to_file(0); \
exit(1); \
} \
} while(0)
/* ===================================================================== */
/* GPU memory management */
/* ===================================================================== */
static void gpu_ensure_staging(int nfloats) {
if (nfloats <= d_staging_n) return;
if (d_staging) CUDA_CHECK(cudaFree(d_staging));
d_staging_n = nfloats;
CUDA_CHECK(cudaMalloc(&d_staging, d_staging_n * sizeof(float)));
}
static void gpu_realloc_if_needed(int required) {
if (required <= g_capacity) return;
int new_cap = g_capacity;
while (new_cap < required) new_cap *= 2;
void *d_new;
CUDA_CHECK(cudaMalloc(&d_new, (size_t)new_cap * g_dim * g_elem_size));
if (d_vectors && g_count > 0) {
CUDA_CHECK(cudaMemcpy(d_new, d_vectors, (size_t)g_count * g_dim * g_elem_size, cudaMemcpyDeviceToDevice));
}
if (d_vectors) CUDA_CHECK(cudaFree(d_vectors));
d_vectors = d_new;
if (d_dists) CUDA_CHECK(cudaFree(d_dists));
CUDA_CHECK(cudaMalloc(&d_dists, new_cap * sizeof(float)));
free(h_dists);
free(h_ids);
h_dists = (float *)malloc(new_cap * sizeof(float));
h_ids = (int *) malloc(new_cap * sizeof(int));
if (!h_dists || !h_ids) {
fprintf(stderr, "FATAL: out of CPU memory allocating distance buffers (%d slots)\n", new_cap);
save_to_file(1);
exit(1);
}
unsigned char *new_alive = (unsigned char *)realloc(g_alive, new_cap);
if (!new_alive) {
fprintf(stderr, "FATAL: out of CPU memory reallocating alive mask (%d bytes)\n", new_cap);
save_to_file(1);
exit(1);
}
memset(new_alive + g_alive_cap, 1, new_cap - g_alive_cap);
g_alive = new_alive;
g_alive_cap = new_cap;
if (d_alive) CUDA_CHECK(cudaFree(d_alive));
CUDA_CHECK(cudaMalloc(&d_alive, new_cap));
CUDA_CHECK(cudaMemcpy(d_alive, g_alive, new_cap, cudaMemcpyHostToDevice));
char **new_labels = (char **)realloc(g_labels, new_cap * sizeof(char *));
if (!new_labels) {
fprintf(stderr, "FATAL: out of CPU memory reallocating label table (%d slots)\n", new_cap);
save_to_file(1);
exit(1);
}
memset(new_labels + g_labels_cap, 0, (new_cap - g_labels_cap) * sizeof(char *));
g_labels = new_labels;
g_labels_cap = new_cap;
unsigned char **new_blobs = (unsigned char **)realloc(g_blobs, new_cap * sizeof(unsigned char *));
if (!new_blobs) {
fprintf(stderr, "FATAL: out of CPU memory reallocating blob table (%d slots)\n", new_cap);
save_to_file(1);
exit(1);
}
memset(new_blobs + g_blobs_cap, 0, (new_cap - g_blobs_cap) * sizeof(unsigned char *));
g_blobs = new_blobs;
unsigned int *new_blob_lens = (unsigned int *)realloc(g_blob_lens, new_cap * sizeof(unsigned int));
if (!new_blob_lens) {
fprintf(stderr, "FATAL: out of CPU memory reallocating blob length table (%d slots)\n", new_cap);
save_to_file(1);
exit(1);
}
memset(new_blob_lens + g_blobs_cap, 0, (new_cap - g_blobs_cap) * sizeof(unsigned int));
g_blob_lens = new_blob_lens;
g_blobs_cap = new_cap;
g_capacity = new_cap;
}
static void gpu_init() {
g_capacity = INITIAL_CAP;
CUDA_CHECK(cudaMalloc(&d_vectors, (size_t)g_capacity * g_dim * g_elem_size));
CUDA_CHECK(cudaMalloc(&d_query, g_dim * g_elem_size));
CUDA_CHECK(cudaMalloc(&d_dists, g_capacity * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_topk_dists, GPU_TOP_K * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_topk_ids, GPU_TOP_K * sizeof(int)));
CUDA_CHECK(cudaMalloc(&d_alive, g_capacity));
CUDA_CHECK(cudaMemset(d_alive, 1, g_capacity));
if (g_fmt == FMT_F16) {
d_staging_n = 1024 * g_dim;
CUDA_CHECK(cudaMalloc(&d_staging, d_staging_n * sizeof(float)));
}
h_dists = (float *)malloc(g_capacity * sizeof(float));
h_ids = (int *) malloc(g_capacity * sizeof(int));
g_alive = (unsigned char *)malloc(g_capacity);
if (!h_dists || !h_ids || !g_alive) {
fprintf(stderr, "FATAL: out of CPU memory during init (%d slots)\n", g_capacity);
exit(1);
}
memset(g_alive, 1, g_capacity);
g_alive_cap = g_capacity;
g_labels = (char **)calloc(g_capacity, sizeof(char *));
if (!g_labels) {
fprintf(stderr, "FATAL: out of CPU memory allocating label table (%d slots)\n", g_capacity);
exit(1);
}
g_labels_cap = g_capacity;
g_blobs = (unsigned char **)calloc(g_capacity, sizeof(unsigned char *));
g_blob_lens = (unsigned int *)calloc(g_capacity, sizeof(unsigned int));
if (!g_blobs || !g_blob_lens) {
fprintf(stderr, "FATAL: out of CPU memory allocating blob tables (%d slots)\n", g_capacity);
exit(1);
}
g_blobs_cap = g_capacity;
}
static void gpu_shutdown() {
if (d_vectors) cudaFree(d_vectors);
if (d_query) cudaFree(d_query);
if (d_dists) cudaFree(d_dists);
if (d_topk_dists) cudaFree(d_topk_dists);
if (d_topk_ids) cudaFree(d_topk_ids);
if (d_alive) cudaFree(d_alive);
if (d_staging) cudaFree(d_staging);
free(h_dists);
free(h_ids);
free(g_alive);
if (g_labels) {
for (int i = 0; i < g_labels_cap; i++) free(g_labels[i]);
free(g_labels);
}
if (g_blobs) {
for (int i = 0; i < g_blobs_cap; i++) free(g_blobs[i]);
free(g_blobs);
}
free(g_blob_lens);
g_hash_index.clear();
}
/* ===================================================================== */
/* Vector operations */
/* ===================================================================== */
static void upload_and_store(const float *h_data, void *d_dest, int nfloats) {
if (g_fmt == FMT_F32) {
CUDA_CHECK(cudaMemcpy(d_dest, h_data, nfloats * sizeof(float), cudaMemcpyHostToDevice));
} else {
gpu_ensure_staging(nfloats);
CUDA_CHECK(cudaMemcpy(d_staging, h_data, nfloats * sizeof(float), cudaMemcpyHostToDevice));
launch_f32_to_f16(d_staging, d_dest, nfloats);
CUDA_CHECK(cudaDeviceSynchronize());
}
}
/* validate label: returns 0 ok, -1 invalid char, -2 too long, -3 empty.
* rejected: control chars, space, and : * ? " < > | ,
* allowed: / \ . _ - = ! ' ^ + % & ( ) [ ] { } @ # $ ~ ` ; alphanumerics, etc. */
static int validate_label(const char *label, int len) {
if (len <= 0) return -3;
if (len > MAX_LABEL_BYTES) return -2;
for (int i = 0; i < len; i++) {
unsigned char c = (unsigned char)label[i];
if (c <= 0x1F || c == 0x7F) return -1;
if (c == ' ' || c == ':' || c == '*' || c == '?' || c == '"' ||
c == '<' || c == '>' || c == '|' || c == ',') return -1;
}
return 0;
}
/* set label without validation — used by .meta load (lenient, matches 1.x) */
static void vec_set_label_raw(int slot, const char *label, int len) {
if (slot < 0 || slot >= g_labels_cap) return;
free(g_labels[slot]);
if (!label || len <= 0) { g_labels[slot] = NULL; return; }
/* strip UTF-8 BOM */
if (len >= 3 && (unsigned char)label[0] == 0xEF &&
(unsigned char)label[1] == 0xBB && (unsigned char)label[2] == 0xBF) {
label += 3; len -= 3;
}
if (len <= 0) { g_labels[slot] = NULL; return; }
char *buf = (char *)malloc(len + 1);
if (!buf) {
fprintf(stderr, "ERROR: out of CPU memory storing label (len=%d), label dropped\n", len);
g_labels[slot] = NULL;
return;
}
memcpy(buf, label, len);
buf[len] = '\0';
g_labels[slot] = buf;
}
/* set label with strict validation — used by all wire write paths (PUSH, CMD_LABEL).
* returns 0 on success, validate_label() error codes otherwise. on error, prior label is unchanged. */
static int vec_set_label(int slot, const char *label, int len) {
if (slot < 0 || slot >= g_labels_cap) return -1;
if (!label || len <= 0) {
free(g_labels[slot]);
g_labels[slot] = NULL;
return 0;
}
int rc = validate_label(label, len);
if (rc != 0) return rc;
char *buf = (char *)malloc(len + 1);
if (!buf) {
fprintf(stderr, "ERROR: out of CPU memory storing label (len=%d), label dropped\n", len);
return -1;
}
memcpy(buf, label, len);
buf[len] = '\0';
free(g_labels[slot]);
g_labels[slot] = buf;
return 0;
}
/* hash the f32 input as it would be stored on the GPU. for f16 DBs we convert
* each lane to IEEE-754 binary16 (round-toward-zero, finite range) and hash the
* resulting unsigned shorts. matches the bytes that GPU upload_and_store produces. */
static uint64_t hash_input_as_stored(const float *h_vec) {
if (g_fmt == FMT_F32) {
return XXH64(h_vec, (size_t)g_dim * sizeof(float), 0);
}
/* f16 path: convert lane-by-lane to binary16. round-toward-zero matches
* the simplest GPU __float2half_rn-equivalent we need for hash agreement;
* exact bit pattern is what matters, the value's accuracy is irrelevant. */
static unsigned short *scratch = NULL;
static int scratch_dim = 0;
if (g_dim > scratch_dim) {
free(scratch);
scratch = (unsigned short *)malloc(g_dim * sizeof(unsigned short));
if (!scratch) return 0;
scratch_dim = g_dim;
}
for (int i = 0; i < g_dim; i++) {
unsigned int f;
memcpy(&f, &h_vec[i], sizeof(unsigned int));
unsigned int sign = (f >> 16) & 0x8000;
int exp32 = (int)((f >> 23) & 0xFF) - 127;
unsigned int mant32 = f & 0x7FFFFF;
unsigned short h;
if (exp32 == 128) {
/* inf/nan */
h = (unsigned short)(sign | 0x7C00 | (mant32 ? 0x0200 : 0));
} else if (exp32 > 15) {
h = (unsigned short)(sign | 0x7C00); /* overflow -> inf */
} else if (exp32 >= -14) {
h = (unsigned short)(sign | (((exp32 + 15) & 0x1F) << 10) | (mant32 >> 13));
} else if (exp32 >= -24) {
unsigned int m = (mant32 | 0x800000) >> (-(exp32 + 14));
h = (unsigned short)(sign | (m >> 13));
} else {
h = (unsigned short)sign; /* underflow -> signed zero */
}
scratch[i] = h;
}
return XXH64(scratch, (size_t)g_dim * sizeof(unsigned short), 0);
}
/* hash the stored vector bytes for an existing slot — reads from GPU. */
static uint64_t hash_slot_stored(int slot) {
static unsigned char *buf = NULL;
static size_t buf_cap = 0;
size_t need = (size_t)g_dim * g_elem_size;
if (need > buf_cap) {
free(buf);
buf = (unsigned char *)malloc(need);
if (!buf) return 0;
buf_cap = need;
}
CUDA_CHECK(cudaMemcpy(buf, (char *)d_vectors + (size_t)slot * g_dim * g_elem_size,
need, cudaMemcpyDeviceToHost));
return XXH64(buf, need, 0);
}
/* erase a (hash, slot) entry from the multimap. no-op if not present. */
static void hash_index_erase(uint64_t hash, uint32_t slot) {
auto range = g_hash_index.equal_range(hash);
for (auto it = range.first; it != range.second; ++it) {
if (it->second == slot) { g_hash_index.erase(it); return; }
}
}
/* find an alive slot whose stored bytes equal h_vec's stored form.
* returns slot index or -1 if no match. caller supplies the input hash. */
static int hash_index_find_match(uint64_t hash, const float *h_vec) {
auto range = g_hash_index.equal_range(hash);
if (range.first == range.second) return -1;
/* compare by byte-exact stored form. read each candidate's GPU bytes
* and memcmp against the same encoding of h_vec. */
static unsigned char *cand = NULL; /* candidate slot bytes */
static unsigned char *want = NULL; /* h_vec encoded as stored */
static size_t cap = 0;
size_t need = (size_t)g_dim * g_elem_size;
if (need > cap) {
free(cand); free(want);
cand = (unsigned char *)malloc(need);
want = (unsigned char *)malloc(need);
if (!cand || !want) { cap = 0; return -1; }
cap = need;
}
/* build "want" buffer: f32 just memcpy, f16 use the same convert as hash_input_as_stored */
if (g_fmt == FMT_F32) {
memcpy(want, h_vec, need);
} else {
unsigned short *w16 = (unsigned short *)want;
for (int i = 0; i < g_dim; i++) {
unsigned int f;
memcpy(&f, &h_vec[i], sizeof(unsigned int));
unsigned int sign = (f >> 16) & 0x8000;
int exp32 = (int)((f >> 23) & 0xFF) - 127;
unsigned int mant32 = f & 0x7FFFFF;
unsigned short h;
if (exp32 == 128) {
h = (unsigned short)(sign | 0x7C00 | (mant32 ? 0x0200 : 0));
} else if (exp32 > 15) {
h = (unsigned short)(sign | 0x7C00);
} else if (exp32 >= -14) {
h = (unsigned short)(sign | (((exp32 + 15) & 0x1F) << 10) | (mant32 >> 13));
} else if (exp32 >= -24) {
unsigned int m = (mant32 | 0x800000) >> (-(exp32 + 14));
h = (unsigned short)(sign | (m >> 13));
} else {
h = (unsigned short)sign;
}
w16[i] = h;
}
}
for (auto it = range.first; it != range.second; ++it) {
uint32_t slot = it->second;
if ((int)slot >= g_count) continue;
if (!g_alive[slot]) continue;
CUDA_CHECK(cudaMemcpy(cand, (char *)d_vectors + (size_t)slot * g_dim * g_elem_size,
need, cudaMemcpyDeviceToHost));
if (memcmp(cand, want, need) == 0) return (int)slot;
}
return -1;
}
/* set blob — opaque bytes, no validation beyond length cap. len=0 clears. */
static int vec_set_blob(int slot, const unsigned char *bytes, unsigned int len) {
if (slot < 0 || slot >= g_blobs_cap) return -1;
if (len > MAX_DATA_BYTES) return -2;
if (!bytes || len == 0) {
free(g_blobs[slot]);
g_blobs[slot] = NULL;
g_blob_lens[slot] = 0;
return 0;
}
unsigned char *buf = (unsigned char *)malloc(len);
if (!buf) {
fprintf(stderr, "ERROR: out of CPU memory storing blob (len=%u)\n", len);
return -1;
}
memcpy(buf, bytes, len);
free(g_blobs[slot]);
g_blobs[slot] = buf;
g_blob_lens[slot] = len;
return 0;
}
static int vec_push(const float *h_vec) {
gpu_realloc_if_needed(g_count + 1);
int slot = g_count;
upload_and_store(h_vec, (char *)d_vectors + (size_t)slot * g_dim * g_elem_size, g_dim);
g_count++;
g_dirty = 1; g_last_write = time(NULL);
g_hash_index.insert({hash_input_as_stored(h_vec), (uint32_t)slot});
return slot;
}
static int vec_pull(const float *h_query, int *out_ids, float *out_dists, int mode) {
int alive = g_count - g_deleted;
if (alive <= 0) return 0;
int n = g_count;
int k = (alive < DEFAULT_TOP_K) ? alive : DEFAULT_TOP_K;
cudaGetLastError();
upload_and_store(h_query, d_query, g_dim);
if (g_fmt == FMT_F32) {
if (mode == 1) launch_cos_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
else launch_l2_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
} else {
if (mode == 1) launch_cos_f16(d_vectors, d_query, d_dists, n, g_dim);
else launch_l2_f16(d_vectors, d_query, d_dists, n, g_dim);
}
CUDA_CHECK(cudaDeviceSynchronize());
if (n >= GPU_SORT_THRESHOLD) {
launch_topk(d_dists, d_alive, n, d_topk_dists, d_topk_ids);
CUDA_CHECK(cudaMemcpy(h_dists, d_topk_dists, k * sizeof(float), cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_ids, d_topk_ids, k * sizeof(int), cudaMemcpyDeviceToHost));
} else {
CUDA_CHECK(cudaMemcpy(h_dists, d_dists, n * sizeof(float), cudaMemcpyDeviceToHost));
for (int i = 0; i < n; i++) {
h_ids[i] = i;
if (!g_alive[i]) h_dists[i] = 3.402823e+38f;
}
for (int i = 0; i < k; i++) {
int best = i;
for (int j = i + 1; j < n; j++) {
if (h_dists[j] < h_dists[best]) best = j;
}
float td = h_dists[i]; h_dists[i] = h_dists[best]; h_dists[best] = td;
int ti = h_ids[i]; h_ids[i] = h_ids[best]; h_ids[best] = ti;
}
}
for (int i = 0; i < k; i++) {
out_ids[i] = h_ids[i];
out_dists[i] = h_dists[i];
}
return k;
}
/* same as vec_pull but the query is an existing slot — copies device-to-device */
static int vec_pull_by_idx(int idx, int *out_ids, float *out_dists, int mode) {
int alive = g_count - g_deleted;
if (alive <= 0) return 0;
int n = g_count;
int k = (alive < DEFAULT_TOP_K) ? alive : DEFAULT_TOP_K;
cudaGetLastError();
CUDA_CHECK(cudaMemcpy(d_query, (char *)d_vectors + (size_t)idx * g_dim * g_elem_size,
g_dim * g_elem_size, cudaMemcpyDeviceToDevice));
if (g_fmt == FMT_F32) {
if (mode == 1) launch_cos_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
else launch_l2_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
} else {
if (mode == 1) launch_cos_f16(d_vectors, d_query, d_dists, n, g_dim);
else launch_l2_f16(d_vectors, d_query, d_dists, n, g_dim);
}
CUDA_CHECK(cudaDeviceSynchronize());
if (n >= GPU_SORT_THRESHOLD) {
launch_topk(d_dists, d_alive, n, d_topk_dists, d_topk_ids);
CUDA_CHECK(cudaMemcpy(h_dists, d_topk_dists, k * sizeof(float), cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_ids, d_topk_ids, k * sizeof(int), cudaMemcpyDeviceToHost));
} else {
CUDA_CHECK(cudaMemcpy(h_dists, d_dists, n * sizeof(float), cudaMemcpyDeviceToHost));
for (int i = 0; i < n; i++) {
h_ids[i] = i;
if (!g_alive[i]) h_dists[i] = 3.402823e+38f;
}
for (int i = 0; i < k; i++) {
int best = i;
for (int j = i + 1; j < n; j++) {
if (h_dists[j] < h_dists[best]) best = j;
}
float td = h_dists[i]; h_dists[i] = h_dists[best]; h_dists[best] = td;
int ti = h_ids[i]; h_ids[i] = h_ids[best]; h_ids[best] = ti;
}
}
for (int i = 0; i < k; i++) {
out_ids[i] = h_ids[i];
out_dists[i] = h_dists[i];
}
return k;
}
/* ===================================================================== */
/* Clustering (DBSCAN) */
/* ===================================================================== */
typedef int (*write_fn)(void *ctx, const char *buf, int len);
#define CLUSTER_UNVISITED -1
#define CLUSTER_NOISE -2
/*
* vec_cluster: DBSCAN clustering using existing GPU distance kernels.
* Each seed does a device-to-device copy into d_query, launches the distance
* kernel, copies distances back, and thresholds to find neighbors.
* No new CUDA kernels needed.
*
* Returns number of clusters found. Writes results via writer callback.
* Format: one line per cluster "id:member,member,...\n", noise last.
*/
static void vec_cluster(float eps, int min_pts, int mode, write_fn writer, void *wctx) {
int n = g_count;
int alive = n - g_deleted;
if (alive <= 0) {
writer(wctx, "end\n", 4);
return;
}
long long t_start = now_ms();
fprintf(stderr, "cluster: %d vectors, eps=%.4g, min_pts=%d, %s\n",
alive, eps, min_pts, mode == 1 ? "cosine" : "L2");
float eps_sq = (mode == 1) ? eps : eps * eps; /* cosine is already 0..2, L2 we store squared */
int *cluster_id = (int *)malloc(n * sizeof(int));
int *queue = (int *)malloc(n * sizeof(int));
float *dists_buf = (float *)malloc(n * sizeof(float));
if (!cluster_id || !queue || !dists_buf) {
fprintf(stderr, "cluster: out of CPU memory (n=%d)\n", n);
free(cluster_id); free(queue); free(dists_buf);
writer(wctx, "err out of memory\n", 18);
return;
}
for (int i = 0; i < n; i++)
cluster_id[i] = g_alive[i] ? CLUSTER_UNVISITED : CLUSTER_NOISE;
int cluster = 0;
for (int i = 0; i < n; i++) {
if (cluster_id[i] != CLUSTER_UNVISITED) continue;
/* compute distances from vector i to all others using GPU */
CUDA_CHECK(cudaMemcpy(d_query, (char *)d_vectors + (size_t)i * g_dim * g_elem_size,
g_dim * g_elem_size, cudaMemcpyDeviceToDevice));
if (g_fmt == FMT_F32) {
if (mode == 1) launch_cos_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
else launch_l2_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
} else {
if (mode == 1) launch_cos_f16(d_vectors, d_query, d_dists, n, g_dim);
else launch_l2_f16(d_vectors, d_query, d_dists, n, g_dim);
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(dists_buf, d_dists, n * sizeof(float), cudaMemcpyDeviceToHost));
/* find neighbors within eps */
int q_head = 0, q_tail = 0;
int neighbor_count = 0;
for (int j = 0; j < n; j++) {
if (!g_alive[j]) continue;
if (dists_buf[j] <= eps_sq) neighbor_count++;
}
if (neighbor_count < min_pts) {
cluster_id[i] = CLUSTER_NOISE;
continue;
}
/* start new cluster — assign seed and enqueue all initial neighbors */
cluster_id[i] = cluster;
for (int j = 0; j < n; j++) {
if (!g_alive[j] || dists_buf[j] > eps_sq) continue;
if (cluster_id[j] == CLUSTER_NOISE) {
cluster_id[j] = cluster;
} else if (cluster_id[j] == CLUSTER_UNVISITED) {
cluster_id[j] = cluster; /* mark immediately to prevent duplicates */
queue[q_tail++] = j;
}
}
/* process queue — expand cluster */
while (q_head < q_tail) {
int j = queue[q_head++];
/* range query from j */
CUDA_CHECK(cudaMemcpy(d_query, (char *)d_vectors + (size_t)j * g_dim * g_elem_size,
g_dim * g_elem_size, cudaMemcpyDeviceToDevice));
if (g_fmt == FMT_F32) {
if (mode == 1) launch_cos_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
else launch_l2_f32((const float *)d_vectors, (const float *)d_query, d_dists, n, g_dim);
} else {
if (mode == 1) launch_cos_f16(d_vectors, d_query, d_dists, n, g_dim);
else launch_l2_f16(d_vectors, d_query, d_dists, n, g_dim);
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(dists_buf, d_dists, n * sizeof(float), cudaMemcpyDeviceToHost));
int j_neighbors = 0;
for (int k = 0; k < n; k++) {
if (!g_alive[k]) continue;
if (dists_buf[k] <= eps_sq) j_neighbors++;
}
if (j_neighbors >= min_pts) {
/* add unvisited neighbors to queue; mark immediately to prevent duplicates */
for (int k = 0; k < n; k++) {
if (!g_alive[k]) continue;
if (dists_buf[k] <= eps_sq) {
if (cluster_id[k] == CLUSTER_NOISE) {
cluster_id[k] = cluster;
} else if (cluster_id[k] == CLUSTER_UNVISITED) {
cluster_id[k] = cluster; /* mark before enqueue — prevents duplicates */
queue[q_tail++] = k;
}
}
}
}
}
cluster++;
}
/* output results */
char line[256];
int line_len;
/* one line per cluster: member,member,...\n */
for (int c = 0; c < cluster; c++) {
int first = 1;
for (int i = 0; i < n; i++) {
if (cluster_id[i] != c) continue;
const char *lbl = (i < g_labels_cap) ? g_labels[i] : NULL;
if (lbl)
line_len = snprintf(line, sizeof(line), "%s%s", first ? "" : ",", lbl);
else
line_len = snprintf(line, sizeof(line), "%s%d", first ? "" : ",", i);
writer(wctx, line, line_len);
first = 0;
}
writer(wctx, "\n", 1);
}
/* noise: same format, one line */
{
int first = 1;
int has_noise = 0;
for (int i = 0; i < n; i++) {
if (cluster_id[i] != CLUSTER_NOISE || !g_alive[i]) continue;
has_noise = 1;
const char *lbl = (i < g_labels_cap) ? g_labels[i] : NULL;
if (lbl)
line_len = snprintf(line, sizeof(line), "%s%s", first ? "" : ",", lbl);
else
line_len = snprintf(line, sizeof(line), "%s%d", first ? "" : ",", i);
writer(wctx, line, line_len);
first = 0;
}
if (has_noise) writer(wctx, "\n", 1);
}
line_len = snprintf(line, sizeof(line), "end\n");
writer(wctx, line, line_len);
/* count noise for logging */
int noise_count = 0;
for (int i = 0; i < n; i++)
if (cluster_id[i] == CLUSTER_NOISE && g_alive[i]) noise_count++;
char el[32]; format_elapsed(now_ms() - t_start, el, sizeof(el));
fprintf(stderr, "cluster: %d clusters, %d noise (%s)\n", cluster, noise_count, el);
free(cluster_id);
free(queue);
free(dists_buf);
}
/* ===================================================================== */
/* Represent: one most-distinct member per DBSCAN cluster */
/* ===================================================================== */
/*
* vec_represent: runs DBSCAN internally, then for each cluster picks the
* member farthest from that cluster's centroid — the most atypical, boundary-
* sitting face rather than the average one. Noise points are skipped.
* Output: one slot index or label per line, then "end\n".
*/
static void vec_represent(float eps, int min_pts, int mode, write_fn writer, void *wctx) {
int n = g_count;
int alive = n - g_deleted;
char resp[256];
int rlen;
if (alive <= 0) { writer(wctx, "end\n", 4); return; }
long long t_start = now_ms();
fprintf(stderr, "represent: %d vectors, eps=%.4g, min_pts=%d, %s\n",
alive, eps, min_pts, mode == 1 ? "cosine" : "L2");
/* ---- phase 1: DBSCAN (identical to vec_cluster) ---- */
int *cluster_id = (int *) malloc(n * sizeof(int));
int *queue = (int *) malloc(n * sizeof(int));
float *dists_buf = (float *)malloc(n * sizeof(float));
if (!cluster_id || !queue || !dists_buf) {
fprintf(stderr, "represent: out of CPU memory (n=%d)\n", n);