Create cub-based argmin primitive and replace argmin_along_rows in ANN kmeans

[Nyrio]

This PR follows up on a suggestion from @cjnolet. The new argmin primitive is up to 5x faster than argmin_along_rows for dimensions relevant to ANN kmeans, and removes code duplication.

The reasons why it is faster are:

• argmin_along_rows often misses on doing a sequential reduction before the tree reduction, especially as it uses large block sizes, as much as 1024.
• CUB has a better reduction algorithm than the basic shared-mem reduction used in argmin_along_rows.
• If we switch the argmin prim to using the cub::BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY algorithm, we can get up to 30% further speedup! (I believe it's safe to use the commutative algorithm here since the offset is contained in the key-value pair so the reduction operation is commutative).

The speedup that I have measured for IVF-Flat build with the InnerProduct metric is around 15%.

[cjnolet]

Thanks for these changes, @Nyrio! Definitely happy to see prims for both argmin and argmax and more consolidation being done on the new ANN algos. Mostly minor things.

[tfeher]

Thanks Louis for the PR, I just have a few small comments.

[Nyrio]

@lowener @cjnolet The new argmax header was slightly wrong, and the test was incorrect as well (see explanation below). Please have a detailed look at my latest change to see if you agree with the way I solved this.

The wrapper takes a row-major matrix view, but the implementation is in column-major semantics. As a result:

• The documentation of the wrapper was incorrect. In row-major semantics, we're finding the column with the minimum index for each row (in column-major it's the other way around).
• The dimension of the output was also incorrect. The wrapper checked that the size of the output was the same as the number of rows of the matrix, but the size of the output is N, which here was the number of columns. Note that I switched the arguments D and N passed to the implementation.
• The test passed mostly by chance, because it created a 4x3 row-major matrix, checking against the expected outputs for a 3x4 row-major matrix, but the errors in the test and wrapper canceled each other. Except that it was checking 4 rows in the devArrMatch despite both the actual and expected output only having 3 rows.

I've kept the column-major convention in the deprecated header, fixed the new ones for the row-major convention, and switched the implementation to drop row and columns in favor of N (number of reductions) and D (elements to reduce) and documented that.

[lowener]

Good catch, your changes on argmax make sense for me!

[Nyrio]

@tfeher and @cjnolet Can you update your review status? Thanks!

[cjnolet]

LGTM. Thanks for fixing the intermediate issues as well!

[tfeher]

Thanks Louis for addressing the issues, it looks good to me!

[cjnolet]

@gpucibot merge