[Bug?] WarpReduce: Unexpected results with logical warp size < 32

Not sure if this is a bug or a feature, but it surely is not the behavior the docs suggest. I tried this with cub 1.8.0 and CUDA 10.1 and 9.2.

Say I have a single block with warpSize*n threads, each threads holds some value thread_data = 1. I set up a warpReduce with a logical warp size of warpSize. Then I perform the sum operation in WarpReduce for each logical warp.
I would expect the aggregate to be warpSize for each first logical lane (threadIdx.x%warpSize == 0).
The docs hint this behavior and it is indeed what happens with a warpSize that is a multiple of 32 (16, 8...). However some warpSize values produces unexpected results. Look at this MWE:

```c++
#include
#include

template
__global__ void kern(){
using WarpReduce = cub::WarpReduce;
__shared__ typename WarpReduce::TempStorage temp_storage[warpsPerBlock];
const auto warp_id = threadIdx.x/warpSize;
const auto lane = threadIdx.x%warpSize;
int thread_data = 1;
int aggregate = WarpReduce(temp_storage[warp_id]).Sum(thread_data);
if(lane == 0){
printf("thread: %d, warp_id: %d, aggregate: %d\n", threadIdx.x, warp_id, aggregate);
}
}

int main(){
constexpr int warpSize = 7;
constexpr int warpsPerBlock = 7;
int nblocks = 1;
kern<<>>();
cudaDeviceSynchronize();
return 0;
}
```
Running it results in the following output:
```bash
$ nvcc -arch=sm_52 mwe.cu -run
thread: 35, warp_id: 5, aggregate: 4
thread: 42, warp_id: 6, aggregate: 4
thread: 0, warp_id: 0, aggregate: 7
thread: 7, warp_id: 1, aggregate: 7
thread: 14, warp_id: 2, aggregate: 7
thread: 21, warp_id: 3, aggregate: 7
thread: 28, warp_id: 4, aggregate: 4
```
The same result is obtained for different architectures. Change warpSize to 8 and you get aggregate = 8 in all lanes = 0. Change it to a value between 9 and 15 and again you get odd behavior past the first few logical warps. Then 16 works fine and the same odd behavior arises until warpSize = 32.

The docs seem to imply that logical_warp_size should not be greater than the real warp size, so anything beyond it I considered expected.
The docs also seem to imply that warpSize does not have to be a whole multiple of 32 by saying:
"[...] Computation is slightly more efficient for [...] The architecture's warp size is a whole multiple of LOGICAL_WARP_THREADS"
under Performance Considerations.