Reduce unrolling factors during tuning

[nicolasvasilache]

For some reason rangeUpTo64 is unsed and unrolling factors are powers of 2 up to 256.
I can correlate that with bad compilation times locally.
Let's use rangeUpTo64 instead which can full unroll by divisors of problem sizes and avoids the 128 and 256 cases. Tested locally, looks good to me.

[skimo-openhub]

For some reason rangeUpTo64 is unsed and unrolling factors are powers of 2 up to 256.

"unused"
Apparently, it used to be used for (only!) b2.

I can correlate that with bad compilation times locally.
Let's use rangeUpTo64 instead which can full unroll by divisors of problem sizes and avoids the 128 and 256 cases. Tested locally, looks good to me.

"fully unroll"?

Why is it important to be able to unroll by divisors?
AFAIU, the unroll factor only specifies an upper bound.

[ftynse]

Apparently, it used to be used for (only!) b2.

I suppose that is the tuner's way of calling blockDim.z, which cannot be greater than 64.

AFAIU, the unroll factor only specifies an upper bound.

That is my understanding too. So divisors will be covered by the closest larger power of two.

[skimo-openhub]

On Thu, Apr 26, 2018 at 09:01:51AM +0000, ftynse wrote: > Apparently, it used to be used for (only!) b2. I suppose that is the tuner's way of calling `blockDim.z`, which cannot be greater than 64.

OK, I didn't know about this restriction. But then why was this limit (apparently) removed from the autotuner? skimo

[ftynse]

But then why was this limit (apparently) removed from the autotuner?

Good question. This may also explain some random "launch out of resources" errors that we saw earlier. @ttheodor ?
But this is not related to the current PR.

[thetheodor]

But then why was this limit (apparently) removed from the autotuner?

My guess is that it got lost in some refactoring.

[nicolasvasilache]

It didn't get lost, it was forcing all grid dimensions to < 64.
@ftynse may be unrelated because of this but we should hardcode it less, can it change on various cards?

[ftynse]

can it change on various cards?

That limitation has been consistent across all compute capabilities so far.

This one however maybe relaxed. Only very old cards have max 512 threads per block; newer ones have 1024 or even 2048. Not sure about 32 as lower bound. We may want some smaller divisors of the size, especially if they are used as tile sizes. OTOH, tightening will kick in in that case and reduce the thread count anyway. Still, there may be a situation where we could have chosen something like 8x128 block size but could not because we don't use less than 32 threads per dimension.

[ftynse]

-Werror kicked in on filterHigherThan

[ftynse]

Addressed this myself