diff --git a/quest/src/core/accelerator.hpp b/quest/src/core/accelerator.hpp index 5a8dc37fb..fc6d002a0 100644 --- a/quest/src/core/accelerator.hpp +++ b/quest/src/core/accelerator.hpp @@ -24,11 +24,36 @@ #include "quest/include/qureg.h" #include "quest/include/matrices.h" +#include "quest/src/core/inliner.hpp" #include "quest/src/core/lists.hpp" #include +/* + * COMPILE-TIME VARIABLES + * + * used by cpu_subroutines.cpp and gpu_subroutines to attemptedly set + * a variable to a value known at compile-time (like a templated function's + * parameter), enabling compile-time optimisations of subsequent code which + * uses the variable such a loop unrolling. If the value is not known at + * compile-time (CompileTimeValue==-1 which indicates a templated function has + * been called with more qubits than it has been explicitly instantiated and + * optimised for), the runtime value is used, precluding optimisations. + * + * TODO: this trick is likely not presently working in (stoopid) MSVC compilers! + */ + +template +INLINE constexpr int accel_tryUseCompileTimeValue(int runtimeValue) { + + if constexpr (CompileTimeValue == -1) + return runtimeValue; + else + return CompileTimeValue; +} + + /* * TEMPLATE INSTANTIATION MACROS * @@ -124,27 +149,6 @@ template returntype funcname <-1,numtargs, conj, haspower> args; - -/* - * COMPILE-TIME VARIABLE MACROS - * - * used by cpu_subroutines.cpp and gpu_subroutines to attemptedly set - * a variable to a value known at compile-time (like a templated function's - * parameter), enabling compile-time optimisations of subsequent code which - * uses the variable such a loop unrolling. If the value is not known at - * compile-time (compileval==-1 which indicates a templated function has - * been called with more qubits than it has been explicitly instantiated and - * optimised for), the runtime value is used, precluding optimisations. - */ - -#define SET_VAR_AT_COMPILE_TIME(type, name, compileval, runtimeval) \ - type name; \ - if constexpr (compileval == -1) \ - name = runtimeval; \ - else \ - name = compileval; - - /* * GETTERS */ diff --git a/quest/src/cpu/cpu_subroutines.cpp b/quest/src/cpu/cpu_subroutines.cpp index 315b41717..c1ba9b3aa 100644 --- a/quest/src/cpu/cpu_subroutines.cpp +++ b/quest/src/cpu/cpu_subroutines.cpp @@ -237,7 +237,7 @@ qindex cpu_statevec_packAmpsIntoBuffer(Qureg qureg, ConstList64 qubitInds, Const auto qubitStateMask = util_getBitMask(qubitInds, qubitStates); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numBits, NumQubits, qubitInds.size()); + int numBits = accel_tryUseCompileTimeValue(qubitInds.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); + int numTargBits = accel_tryUseCompileTimeValue(targs.size()); int numQubitBits = numCtrlBits + numTargBits; - qindex numTargAmps = powerOf2(numTargBits); + auto numTargAmps = powerOf2(numTargBits); // create an explicit parallel region to avoid re-initialisation of vectors every iteration #pragma omp parallel if(qureg.isMultithreaded) @@ -678,7 +672,7 @@ void cpu_statevec_anyCtrlOneTargDiagMatr_sub(Qureg qureg, ConstList64 ctrls, Con auto ctrlStateMask = util_getBitMask(ctrls, ctrlStates); // use template params to compile-time unroll loops in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); + int numTargBits = accel_tryUseCompileTimeValue(targs.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(numXY); // iA = nth local index where targs have value m, iB = (last - nth) such index qindex iA = setBits(i0, indXY, numTargBits, v); @@ -1004,12 +998,12 @@ void cpu_statevector_anyCtrlPauliTensorOrGadget_subA( auto maskYZ = util_getBitMask(util_getConcatenated(y, z)); // use template params to compile-time unroll loops in insertBits() and inner-loop below - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); - SET_VAR_AT_COMPILE_TIME(int, numTargBits, NumTargs, sortedTargsXY.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); + int numTargBits = accel_tryUseCompileTimeValue(sortedTargsXY.size()); int numQubitBits = numCtrlBits + numTargBits; // each outer iteration handles all assignments of the target qubits, and each ctrl halves the outer iterations - qindex numOuterIts = qureg.numAmpsPerNode / powerOf2(numCtrlBits + numTargBits); + qindex numOuterIts = qureg.numAmpsPerNode / powerOf2(numQubitBits); // each inner iteration modifies 2 amplitudes (may be compile-time sized) qindex numInnerIts = powerOf2(numTargBits) / 2; // divides evenly @@ -1088,7 +1082,7 @@ void cpu_statevector_anyCtrlPauliTensorOrGadget_subB( auto maskYZ = util_getBitMask(util_getConcatenated(y, z)); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n(ctrls.size()); #pragma omp parallel for if(qureg.isMultithreaded) for (qindex n=0; n coeffs qindex numIts = outQureg.numAmpsPerNode; // use template param to compile-time unroll inner loop below - SET_VAR_AT_COMPILE_TIME(int, numQuregs, NumQuregs, inQuregs.size()); + int numQuregs = accel_tryUseCompileTimeValue(inQuregs.size()); #pragma omp parallel for if(outQureg.isMultithreaded) for (qindex n=0; n(targs.size()); // may be inferred at compile-time - int numAllTargs = 2*numTargPairs; + int numAllTargs = 2 * numTargPairs; qindex numInnerIts = powerOf2(numTargPairs); /// @todo @@ -2027,7 +2021,7 @@ qreal cpu_statevec_calcProbOfMultiQubitOutcome_sub(Qureg qureg, ConstList64 qubi auto qubitStateMask = util_getBitMask(qubits, outcomes); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numBits, NumQubits, qubits.size()); + int numBits = accel_tryUseCompileTimeValue(qubits.size()); #pragma omp parallel for reduction(+:prob) if(qureg.isMultithreaded) for (qindex n=0; n(qubits.size()); #pragma omp parallel for reduction(+:prob) if(qureg.isMultithreaded) for (qindex n=0; n(qubits.size()); qindex numOutcomes = powerOf2(numBits); // decide whether to parallelise below amp-clearing, since outProbs ~ dim of a qureg @@ -2146,7 +2140,7 @@ void cpu_densmatr_calcProbsOfAllMultiQubitOutcomes_sub(qreal* outProbs, Qureg qu const bool areQubitsSorted = util_isSorted(qubits); // use template param to compile-time unroll loop in getValueOfBits() - SET_VAR_AT_COMPILE_TIME(int, numBits, NumQubits, qubits.size()); + int numBits = accel_tryUseCompileTimeValue(qubits.size()); qindex numOutcomes = powerOf2(numBits); // decide whether to parallelise below amp-clearing, since outProbs ~ dim of a qureg diff --git a/quest/src/gpu/gpu_kernels.cuh b/quest/src/gpu/gpu_kernels.cuh index 65f277483..42617c4a6 100644 --- a/quest/src/gpu/gpu_kernels.cuh +++ b/quest/src/gpu/gpu_kernels.cuh @@ -24,6 +24,7 @@ #include "quest/src/core/bitwise.hpp" #include "quest/src/core/fastmath.hpp" +#include "quest/src/core/accelerator.hpp" #include "quest/src/gpu/gpu_qcomp.cuh" #if ! QUEST_COMPILE_CUDA @@ -100,7 +101,7 @@ __forceinline__ __device__ int cudaGetBitMaskParity(qindex mask) { */ -template +template __global__ void kernel_statevec_packAmpsIntoBuffer( gpu_qcomp* amps, gpu_qcomp* buffer, qindex numThreads, _GRID_CONST_OPT const List64 qubits, qindex qubitStateMask @@ -108,7 +109,7 @@ __global__ void kernel_statevec_packAmpsIntoBuffer( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numBits, NumCtrls, qubits.size()); + int numBits = accel_tryUseCompileTimeValue(qubits.size()); // i = nth local index where qubits are active qindex i = insertBitsWithMaskedValues(n, qubits.data(), numBits, qubitStateMask); @@ -147,12 +148,10 @@ __global__ void kernel_statevec_anyCtrlSwap_subA( ) { GET_THREAD_IND(n, numThreads); - // beware ctrlsAndTargs contains the two targets - constexpr int numTargs = 2; - // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTargs.size() - numTargs); - int numQubitBits = numCtrlBits + numTargs; + int numTargs = 2; + int numCtrls = ctrlsAndTargs.size() - numTargs; + int numQubitBits = numTargs + accel_tryUseCompileTimeValue(numCtrls); // i01 = nth local index where ctrls are active, targ2=0 and targ1=1 qindex i01 = insertBitsWithMaskedValues(n, ctrlsAndTargs.data(), numQubitBits, ctrlsAndTargsMask); @@ -173,7 +172,7 @@ __global__ void kernel_statevec_anyCtrlSwap_subB( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); // i = nth local index where ctrls are active qindex i = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -190,12 +189,10 @@ __global__ void kernel_statevec_anyCtrlSwap_subC( ) { GET_THREAD_IND(n, numThreads); - // beware ctrlsAndTarg contains the single target - constexpr int numTargs = 1; - // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTarg.size() - numTargs); - int numQubitBits = numCtrlBits + numTargs; + int numTargs = 1; + int numCtrls = ctrlsAndTarg.size() - numTargs; + int numQubitBits = numTargs + accel_tryUseCompileTimeValue(numCtrls); // i = nth local index where ctrls and targ are in specified states qindex i = insertBitsWithMaskedValues(n, ctrlsAndTarg.data(), numQubitBits, ctrlsAndTargMask); @@ -220,12 +217,10 @@ __global__ void kernel_statevec_anyCtrlOneTargDenseMatr_subA( ) { GET_THREAD_IND(n, numThreads); - // beware ctrlsAndTarg contains the single target - constexpr int numTargs = 1; - // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTarg.size() - numTargs); - int numQubitBits = numCtrlBits + numTargs; + int numTargs = 1; + int numCtrls = ctrlsAndTarg.size() - numTargs; + int numQubitBits = numTargs + accel_tryUseCompileTimeValue(numCtrls); // i0 = nth local index where ctrls are active and targ is 0 qindex i0 = insertBitsWithMaskedValues(n, ctrlsAndTarg.data(), numQubitBits, ctrlStateMask); @@ -249,7 +244,7 @@ __global__ void kernel_statevec_anyCtrlOneTargDenseMatr_subB( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); // i = nth local index where ctrl bits are active qindex i = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -277,12 +272,10 @@ __global__ void kernel_statevec_anyCtrlTwoTargDenseMatr_sub( ) { GET_THREAD_IND(n, numThreads); - // beware ctrlsAndTargs contains the two targets - constexpr int numTargs = 2; - // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTargs.size() - numTargs); - int numQubitBits = numCtrlBits + numTargs; + int numTargs = 2; + int numCtrls = ctrlsAndTargs.size() - numTargs; + int numQubitBits = numTargs + accel_tryUseCompileTimeValue(numCtrls); // i00 = nth local index where ctrls are active and both targs are 0 qindex i00 = insertBitsWithMaskedValues(n, ctrlsAndTargs.data(), numQubitBits, ctrlStateMask); @@ -327,31 +320,36 @@ __global__ void kernel_statevec_anyCtrlFewTargDenseMatr( ) { GET_THREAD_IND(n, numThreads); - // it is gauranteed that NumTargs <= 5, such that the thread-private array is - // <= 2^5 = 32 amps <= 512 bytes, which aggregated between all threads in the - // block (assumed ~128) is 64 KiB, and which should be small enough to fit into - // SM registers without spillage into slow local memory. This is despite it - // exceeding the maximum per-block shared memory of 48 KiB. Access to this cache - // must be strictly through compile-time-known indices, otherwise it will auto- - // spill to local memory). Hence, this _subA() function is not a subroutine - // despite some logic being common to non-compile-time _subB(), and hence + // it is guaranteed that numTargs <= 5, so we know NumTargAmps <= 32 + static_assert(NumTargs != -1); + static_assert(NumTargs <= 5); + constexpr qindex NumTargAmps = 1 << NumTargs; // = powerOf2(NumTargs), which is not constexpr + + // Since NumTargAmps <= 32, the thread-private array below is <= 512 bytes big, + // which when aggregated between all threads in the block (assumed ~128) is 64 KiB, + // and which should be small enough to fit into SM registers without spillage into + // slow local memory. This is despite it exceeding the maximum per-block shared memory + // of 48 KiB. Access to this cache must be strictly through compile-time-known indices, + // otherwise it will auto-spill to local memory. Hence, this _subA() function is not a + // subroutine despite some logic being common to non-compile-time _subB(), and hence // why the loops below are explicitly compile-time unrolled. Beware that when // numThreadsPerBlock is increased from 128, this kernel will still behave // correctly, but privateCache below will spill over into local memory at a // performance penalty for NumTargs <= 5, with spillage occurring for fewer // NumTargs as numThreadsPerBlock increases. - REGISTER gpu_qcomp privateCache[1 << NumTargs]; + REGISTER gpu_qcomp privateCache[NumTargAmps]; - // we know NumTargs <= 5, though NumCtrls is permitted anything (including -1) - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTargs.size() - targs.size()); - constexpr qindex numTargAmps = (1 << NumTargs); // explicit, in lieu of powerOf2 + // we know NumTargs <= 5, though NumCtrls is permitted anything (including -1), so + // we can only ATTEMPT to compile-time unroll the loop inside insertBits() + int numCtrls = ctrlsAndTargs.size() - targs.size(); + int numQubitBits = NumTargs + accel_tryUseCompileTimeValue(numCtrls); // i0 = nth local index where ctrls are active and targs are all zero - qindex i0 = insertBitsWithMaskedValues(n, ctrlsAndTargs.data(), numCtrlBits + NumTargs, ctrlsAndTargsMask); // loop may be unrolled + qindex i0 = insertBitsWithMaskedValues(n, ctrlsAndTargs.data(), numQubitBits, ctrlsAndTargsMask); // loop may be unrolled // populate cache (force unroll to ensure compile-time cache indices) #pragma unroll - for (qindex k=0; k5 is always unknown/runtime - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTargs.size() - targs.size()); + // there is actually NO POINT to attempt to use the compile-time NumCtrls value + // here, since NumCtrls will be summed with the always-runtime targs.size() + // variable, and this function is the "many-targ" scenario (i.e. numTargs>5); + // but we do it anyway just for consistency + int numTargs = targs.size(); + int numCtrls = ctrlsAndTargs.size() - numTargs; + int numQubits = numTargs + accel_tryUseCompileTimeValue(numCtrls); // always run-time! // unlike all other kernels, each thread modifies multiple batches of amplitudes for (qindex b=0; b(ctrls.size()); // j = nth local index where ctrls are active (in the specified states) qindex j = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -523,7 +526,7 @@ __global__ void kernel_statevec_anyCtrlTwoTargDiagMatr_sub( /// We should verify this! // use template params to compile-time unroll loops in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); // j = nth local index where ctrls are active (in the specified states) qindex j = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -565,8 +568,8 @@ __global__ void kernel_statevec_anyCtrlAnyTargDiagMatr_sub( /// We should verify this! // use template params to compile-time unroll loops in insertBits() and getValueOfBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); - SET_VAR_AT_COMPILE_TIME(int, numTargBits, NumTargs, targs.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); + int numTargBits = accel_tryUseCompileTimeValue(targs.size()); // j = nth local index where ctrls are active (in the specified states) qindex j = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -650,8 +653,9 @@ __global__ void kernel_statevector_anyCtrlPauliTensorOrGadget_subA( GET_THREAD_IND(t, numThreads); // use template params to compile-time unroll loops in insertBits() and setBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrlsAndTargs.size() - targsXY.size()); - SET_VAR_AT_COMPILE_TIME(int, numTargBits, NumTargs, targsXY.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrlsAndTargs.size() - targsXY.size()); + int numTargBits = accel_tryUseCompileTimeValue(targsXY.size()); + int numBothBits = numCtrlBits + numTargBits; // n = local index of amp sub-batch with common i0, v = value of target bits qindex numInnerIts = powerOf2(numTargBits) / 2; @@ -659,7 +663,7 @@ __global__ void kernel_statevector_anyCtrlPauliTensorOrGadget_subA( qindex v = t % numInnerIts; // i0 = nth local index where ctrls are active and targs are all zero (loop therein may be unrolled) - qindex i0 = insertBitsWithMaskedValues(n, ctrlsAndTargs.data(), numCtrlBits + numTargBits, ctrlsAndTargsStateMask); + qindex i0 = insertBitsWithMaskedValues(n, ctrlsAndTargs.data(), numBothBits, ctrlsAndTargsStateMask); // iA = nth local index where targs have value v, iB = (last - nth) such index qindex iA = setBits(i0, targsXY.data(), numTargBits, v); // may be unrolled @@ -690,7 +694,7 @@ __global__ void kernel_statevector_anyCtrlPauliTensorOrGadget_subB( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); // i = nth local index where ctrl bits are in specified states qindex i = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -725,7 +729,7 @@ __global__ void kernel_statevector_anyCtrlAnyTargZOrPhaseGadget_sub( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numCtrlBits, NumCtrls, ctrls.size()); + int numCtrlBits = accel_tryUseCompileTimeValue(ctrls.size()); // i = nth local index where ctrl bits are in specified states qindex i = insertBitsWithMaskedValues(n, ctrls.data(), numCtrlBits, ctrlStateMask); @@ -752,7 +756,7 @@ __global__ void kernel_statevec_setQuregToWeightedSum_sub( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll below loop - SET_VAR_AT_COMPILE_TIME(int, numInner, NumQuregs, numQuregs); + int numInner = accel_tryUseCompileTimeValue(numQuregs); gpu_qcomp amp = getGpuQcomp(0, 0); @@ -1174,7 +1178,7 @@ __global__ void kernel_densmatr_partialTrace_sub( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll below loops - SET_VAR_AT_COMPILE_TIME(int, numTargPairs, NumTargs, ketTargs.size()); + int numTargPairs = accel_tryUseCompileTimeValue(ketTargs.size()); // may be inferred at compile-time int numAllTargs = 2 * numTargPairs; @@ -1227,7 +1231,7 @@ __global__ void kernel_statevec_calcProbsOfAllMultiQubitOutcomes_sub( /// whether this is worthwhile and faster! // use template param to compile-time unroll below loops - SET_VAR_AT_COMPILE_TIME(int, numBits, NumQubits, qubits.size()); + int numBits = accel_tryUseCompileTimeValue(qubits.size()); qreal prob = norm(amps[n]); @@ -1251,7 +1255,7 @@ __global__ void kernel_densmatr_calcProbsOfAllMultiQubitOutcomes_sub( GET_THREAD_IND(n, numThreads); // use template param to compile-time unroll loop in insertBits() - SET_VAR_AT_COMPILE_TIME(int, numBits, NumQubits, qubits.size()); + int numBits = accel_tryUseCompileTimeValue(qubits.size()); // i = index of nth local diagonal elem qindex i = fast_getQuregLocalIndexOfDiagonalAmp(n, firstDiagInd, numAmpsPerCol); diff --git a/quest/src/gpu/gpu_thrust.cuh b/quest/src/gpu/gpu_thrust.cuh index b6c113e73..ae226545c 100644 --- a/quest/src/gpu/gpu_thrust.cuh +++ b/quest/src/gpu/gpu_thrust.cuh @@ -42,6 +42,7 @@ #include "quest/src/core/utilities.hpp" #include "quest/src/core/randomiser.hpp" #include "quest/src/core/fastmath.hpp" +#include "quest/src/core/accelerator.hpp" #include "quest/src/comm/comm_config.hpp" #include @@ -450,7 +451,7 @@ struct functor_insertBits { __host__ __device__ qindex operator()(qindex i) { // use the compile-time value if possible, to auto-unroll the insertBits loop - SET_VAR_AT_COMPILE_TIME(int, nbits, NumBits, numBits); + int nbits = accel_tryUseCompileTimeValue(numBits); // return ith local index where bits have the specified values at the specified indices return insertBitsWithMaskedValues(i, sortedInds.data(), nbits, valueMask);