Skip to content

Commit fe446e4

Browse files
mdouzemdouze
committed
bugfixes
1 parent 250a3d3 commit fe446e4

21 files changed

Lines changed: 243 additions & 154 deletions

‎AutoTune.cpp‎

Lines changed: 13 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -431,7 +431,8 @@ void ParameterSpace::set_index_parameter (
431431
// and fall through to also enable it on sub-indexes
432432
}
433433
if (DC (IndexPreTransform)) {
434-
index = ix->index;
434+
set_index_parameter (ix->index, name, val);
435+
return;
435436
}
436437
if (DC (IndexShards)) {
437438
// call on all sub-indexes
@@ -440,30 +441,28 @@ void ParameterSpace::set_index_parameter (
440441
}
441442
return;
442443
}
443-
if (name == "verbose") {
444-
index->verbose = int(val);
445-
// in case it was an IndexPreTransform
446-
}
447444
if (DC (IndexRefineFlat)) {
448445
if (name == "k_factor_rf") {
449446
ix->k_factor = int(val);
450447
return;
451448
}
452-
index = ix->base_index;
453-
}
454-
if (DC (IndexPreTransform)) {
455-
index = ix->index;
449+
// otherwise it is for the sub-index
450+
set_index_parameter (&ix->refine_index, name, val);
451+
return;
456452
}
453+
457454
if (name == "verbose") {
458455
index->verbose = int(val);
459456
return; // last verbose that we could find
460457
}
458+
461459
if (name == "nprobe") {
462460
if ( DC(IndexIVF)) {
463461
ix->nprobe = int(val);
464462
return;
465463
}
466464
}
465+
467466
if (name == "ht") {
468467
if (DC (IndexPQ)) {
469468
if (val >= ix->pq.code_size * 8) {
@@ -685,7 +684,6 @@ Index *index_factory (int d, const char *description_in, MetricType metric)
685684
tok;
686685
tok = strtok_r (nullptr, " ,", &ptr)) {
687686
int d_out, opq_M, nbit, M, M2;
688-
char option[100];
689687
std::string stok(tok);
690688

691689
// to avoid mem leaks with exceptions:
@@ -776,10 +774,9 @@ Index *index_factory (int d, const char *description_in, MetricType metric)
776774
del_coarse_quantizer.release ();
777775
index_ivf->own_fields = true;
778776
index_1 = index_ivf;
779-
} else if (!index && sscanf (tok, "PQ%d%10s", &M, option) == 2) {
780-
std::string soption = option;
781-
// np to disable polysemous trainign
782-
FAISS_THROW_IF_NOT(soption == "" || soption == "np");
777+
} else if (!index && (sscanf (tok, "PQ%d", &M) == 1 ||
778+
sscanf (tok, "PQ%dnp", &M) == 1)) {
779+
bool do_polysemous_training = stok.find("np") == std::string::npos;
783780
if (coarse_quantizer) {
784781
IndexIVFPQ *index_ivf = new IndexIVFPQ (
785782
coarse_quantizer, d, ncentroids, M, 8);
@@ -789,11 +786,11 @@ Index *index_factory (int d, const char *description_in, MetricType metric)
789786
index_ivf->cp.spherical = metric == METRIC_INNER_PRODUCT;
790787
del_coarse_quantizer.release ();
791788
index_ivf->own_fields = true;
792-
index_ivf->do_polysemous_training = soption != "np";
789+
index_ivf->do_polysemous_training = do_polysemous_training;
793790
index_1 = index_ivf;
794791
} else {
795792
IndexPQ *index_pq = new IndexPQ (d, M, 8, metric);
796-
index_pq->do_polysemous_training = soption != "np";
793+
index_pq->do_polysemous_training = do_polysemous_training;
797794
index_1 = index_pq;
798795
}
799796

‎IndexPQ.cpp‎

Lines changed: 15 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -260,7 +260,6 @@ static size_t polysemous_inner_loop (
260260
void IndexPQ::search_core_polysemous (idx_t n, const float *x, idx_t k,
261261
float *distances, idx_t *labels) const
262262
{
263-
FAISS_THROW_IF_NOT (pq.code_size % 8 == 0);
264263
FAISS_THROW_IF_NOT (pq.byte_per_idx == 1);
265264

266265
// PQ distance tables
@@ -319,12 +318,17 @@ void IndexPQ::search_core_polysemous (idx_t n, const float *x, idx_t k,
319318
(*this, dis_table_qi, q_code, k, heap_dis, heap_ids);
320319
break;
321320
default:
322-
if (pq.code_size % 8 == 0)
321+
if (pq.code_size % 8 == 0) {
323322
n_pass += polysemous_inner_loop<HammingComputerM8>
324323
(*this, dis_table_qi, q_code, k, heap_dis, heap_ids);
325-
else
324+
} else if (pq.code_size % 4 == 0) {
326325
n_pass += polysemous_inner_loop<HammingComputerM4>
327326
(*this, dis_table_qi, q_code, k, heap_dis, heap_ids);
327+
} else {
328+
FAISS_THROW_FMT(
329+
"code size %zd not supported for polysemous",
330+
pq.code_size);
331+
}
328332
break;
329333
}
330334
} else {
@@ -342,8 +346,14 @@ void IndexPQ::search_core_polysemous (idx_t n, const float *x, idx_t k,
342346
(*this, dis_table_qi, q_code, k, heap_dis, heap_ids);
343347
break;
344348
default:
345-
n_pass += polysemous_inner_loop<GenHammingComputerM8>
346-
(*this, dis_table_qi, q_code, k, heap_dis, heap_ids);
349+
if (pq.code_size % 8 == 0) {
350+
n_pass += polysemous_inner_loop<GenHammingComputerM8>
351+
(*this, dis_table_qi, q_code, k, heap_dis, heap_ids);
352+
} else {
353+
FAISS_THROW_FMT(
354+
"code size %zd not supported for polysemous",
355+
pq.code_size);
356+
}
347357
break;
348358
}
349359
}

‎gpu/GpuAutoTune.cpp‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -381,20 +381,20 @@ void GpuParameterSpace::set_index_parameter (
381381
set_index_parameter (ix->at(i), name, val);
382382
return;
383383
}
384-
if (DC (GpuIndexIVF)) {
385-
if (name == "nprobe") {
384+
if (name == "nprobe") {
385+
if (DC (GpuIndexIVF)) {
386386
ix->setNumProbes (int (val));
387387
return;
388388
}
389389
}
390-
if(DC (GpuIndexIVFPQ)) {
391-
if (name == "use_precomputed_table") {
390+
if (name == "use_precomputed_table") {
391+
if (DC (GpuIndexIVFPQ)) {
392392
ix->setPrecomputedCodes(bool (val));
393393
return;
394394
}
395395
}
396396

397-
// maybe norma lindex parameters apply?
397+
// maybe normal index parameters apply?
398398
ParameterSpace::set_index_parameter (index, name, val);
399399
}
400400

‎gpu/GpuIndexIVF.cu‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,9 @@ GpuIndexIVF::GpuIndexIVF(GpuResources* resources,
2929
nprobe_(1),
3030
quantizer_(nullptr) {
3131
#ifndef FAISS_USE_FLOAT16
32-
FAISS_THROW_IF_NOT_MSG(!ivfConfig_.flatConfig.useFloat16CoarseQuantizer,
33-
"float16 unsupported; need CUDA SDK >= 7.5");
32+
FAISS_THROW_IF_NOT_MSG(!ivfConfig_.flatConfig.useFloat16 &&
33+
!ivfConfig_.flatConfig.useFloat16Accumulator,
34+
"float16 unsupported; need CUDA SDK >= 7.5");
3435
#endif
3536

3637
init_();

‎gpu/GpuIndexIVFPQ.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -60,7 +60,7 @@ GpuIndexIVFPQ::GpuIndexIVFPQ(GpuResources* resources,
6060
reserveMemoryVecs_(0),
6161
index_(nullptr) {
6262
#ifndef FAISS_USE_FLOAT16
63-
FAISS_ASSERT(!useFloat16LookupTables_);
63+
FAISS_ASSERT(!config.useFloat16LookupTables);
6464
#endif
6565

6666
verifySettings_();

‎gpu/impl/IVFFlat.cu‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -45,10 +45,6 @@ IVFFlat::IVFFlat(GpuResources* resources,
4545
space),
4646
l2Distance_(l2Distance),
4747
useFloat16_(useFloat16) {
48-
#ifndef FAISS_USE_FLOAT16
49-
FAISS_ASSERT_MSG(!useFloat16, "float16 unsupported");
50-
useFloat16_ = false;
51-
#endif
5248
}
5349

5450
IVFFlat::~IVFFlat() {
@@ -95,6 +91,9 @@ IVFFlat::addCodeVectorsFromCpu(int listId,
9591
lengthInBytes,
9692
stream,
9793
true /* exact reserved size */);
94+
#else
95+
// we are not compiling with float16 support
96+
FAISS_ASSERT(false);
9897
#endif
9998
} else {
10099
listData->append((unsigned char*) vecs,

‎gpu/impl/PQScanMultiPassNoPrecomputed.cu‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -506,7 +506,6 @@ void runPQScanMultiPassNoPrecomputed(Tensor<float, 2, true>& queries,
506506
}
507507
#else
508508
FAISS_ASSERT(!useFloat16Lookup);
509-
int codeSize = sizeof(float);
510509
#endif
511510

512511
int totalCodeDistancesSize =

‎gpu/test/test_gpu_index.py‎

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -7,16 +7,12 @@
77
#! /usr/bin/env python2
88

99
import time
10-
import libfb.py.mkl # noqa
11-
10+
import unittest
1211
import numpy as np
13-
14-
from libfb import testutil
15-
1612
import faiss
1713

1814

19-
class EvalIVFPQAccuracy(testutil.BaseFacebookTestCase):
15+
class EvalIVFPQAccuracy(unittest.TestCase):
2016

2117
def get_dataset(self, small_one=False):
2218
if not small_one:
@@ -110,3 +106,9 @@ def test_cpu_to_gpu_IVFPQ(self):
110106

111107
def test_cpu_to_gpu_IVFFlat(self):
112108
self.do_cpu_to_gpu('IVF128,Flat')
109+
110+
def test_set_gpu_param(self):
111+
index = faiss.index_factory(12, "PCAR8,IVF10,PQ4")
112+
res = faiss.StandardGpuResources()
113+
gpu_index = faiss.index_cpu_to_gpu(res, 0, index)
114+
faiss.GpuParameterSpace().set_index_parameter(index, "nprobe", 3)

‎gpu/utils/DeviceDefs.cuh‎

Lines changed: 7 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@
1313
namespace faiss { namespace gpu {
1414

1515
#ifdef __CUDA_ARCH__
16-
#if __CUDA_ARCH__ <= 620
16+
#if __CUDA_ARCH__ <= 700
1717
constexpr int kWarpSize = 32;
1818
#else
1919
#error Unknown __CUDA_ARCH__; please define parameters for compute capability
@@ -25,37 +25,15 @@ constexpr int kWarpSize = 32;
2525
constexpr int kWarpSize = 32;
2626
#endif // !__CUDA_ARCH__
2727

28+
// This is a memory barrier for intra-warp writes to shared memory.
2829
__forceinline__ __device__ void warpFence() {
29-
// Technically, memory barriers are required via the CUDA
30-
// programming model, since warp synchronous programming no longer
31-
// is guaranteed.
32-
//
33-
// There are two components to it:
34-
// -a barrier known to the compiler such that the compiler will not
35-
// schedule loads and stores across the barrier;
36-
// -a HW-level barrier that guarantees that writes are seen in the
37-
// proper order
38-
//
39-
// However, __threadfence_block() is a stronger constraint than what
40-
// we really want out of the hardware: a warp-wide barrier.
41-
//
42-
// In current hardware, it appears that warp synchronous programming
43-
// is a reality; by all tests it appears safe and race-free.
44-
//
45-
// However, understandably it may not be in the future (based on
46-
// what Nvidia says in the Kepler guide, it may change depending
47-
// upon compiler/toolchain issues or future hardware).
48-
//
49-
// Removing the fence results in 10%+ faster performance.
50-
// However, we are judicious as to where we insert the fence, so if
51-
// this reality ever changes, uncommenting this will result in CUDA
52-
// programming model-safe ordering again.
53-
//
54-
// FIXME: we should probably qualify as volatile as well, since the
55-
// compiler could technically preserve values across loops? This
56-
// seems very impractical for the compiler to do, however.
5730

31+
#if __CUDA_ARCH__ >= 700
32+
__syncwarp();
33+
#else
34+
// For the time being, assume synchronicity.
5835
// __threadfence_block();
36+
#endif
5937
}
6038

6139
} } // namespace

‎gpu/utils/Float16.cu‎

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -48,10 +48,16 @@ void runConvertToFloat32(float* out,
4848
in, in + num, out, HalfToFloat());
4949
}
5050

51-
half hostFloat2Half(float a) {
52-
half h;
51+
__half hostFloat2Half(float a) {
52+
#if CUDA_VERSION >= 9000
53+
__half_raw raw;
54+
raw.x = cpu_float2half_rn(a).x;
55+
return __half(raw);
56+
#else
57+
__half h;
5358
h.x = cpu_float2half_rn(a).x;
5459
return h;
60+
#endif
5561
}
5662

5763
} } // namespace

0 commit comments

Comments
 (0)