`CachingDeviceAllocator` : Reusing allocations is not safe with per-thread implicit streams

After DeviceFree, an allocation can in theory be reused on the same stream it was previously allocated. stream-ordering ensures that all operations before DeviceFree are completed before accessing the reused allocation.

This is implemented as follows
```
// To prevent races with reusing blocks returned by the host but still
// in use by the device, only consider cached blocks that are
// either (from the active stream) or (from an idle stream)
if ((active_stream == block_itr->associated_stream) ||
(CubDebug(cudaEventQuery(block_itr->ready_event) != cudaErrorNotReady)))
{
```

However, this is not sufficient for cudaStreamPerThread, or default stream compiled with `--default-stream per-thread`. These streams have the same numeric value on each thread, respectively, but represent a different stream on each thread. This can lead to invalid sharing of pointers. The following code demonstrates the issue

```
// nvcc -arch=sm_61 -O3 main.cu -o main

#include

#include
#include
#include
#include

__global__
void computekernel(float* data, int N, int iterations){
for(int i = 0; i < N; i++){
for(int iter = 0; iter < iterations; iter++){
data[i] = sin(cos(data[i]));
}
}
}

__global__
void evilzerokernel(float* data, int N){
const int tid = threadIdx.x + blockIdx.x * blockDim.x;
const int stride = blockDim.x * gridDim.x;

for(int i = tid; i < N - 1; i += stride){
data[i] = 0;
}
}

__global__
void checkkernel(const float* data, int N){
const int tid = threadIdx.x + blockIdx.x * blockDim.x;
const int stride = blockDim.x * gridDim.x;

//check that right neighbor is equal to current
for(int i = tid; i < N - 1; i += stride){
bool equal = data[i] == data[i+1];
if(!equal){
printf("error position %d, %f %f\n", i, data[i], data[i+1]);
}
}
}

int main(){

const int N = 1024 * 1024 * 1;

constexpr unsigned int cub_CachingDeviceAllocator_INVALID_BIN = (unsigned int) -1;
constexpr size_t cub_CachingDeviceAllocator_INVALID_SIZE = (size_t) -1;

cub::CachingDeviceAllocator allocator(
2, //bin_growth
1, //min_bin
cub_CachingDeviceAllocator_INVALID_BIN, //max_bin
cub_CachingDeviceAllocator_INVALID_SIZE, //max_cached_bytes
false, //skip_cleanup
true //debug
);

#if 1
cudaStream_t stream1 = cudaStreamPerThread;
cudaStream_t stream2 = cudaStreamPerThread;
#else
cudaStream_t stream1; cudaStreamCreate(&stream1);
cudaStream_t stream2; cudaStreamCreate(&stream2);
#endif

float* d_data1;
allocator.DeviceAllocate((void**)&d_data1, sizeof(float) * N, stream1);

cudaMemsetAsync(d_data1, 42, sizeof(float) * N, stream1);

computekernel<<<1,1,0,stream1>>>(d_data1, N, 32); //change iteration count (32) such that the second cpu thread finishes before checkkernel

checkkernel<<<1024, 256, 0, stream1>>>(d_data1, N);

allocator.DeviceFree(d_data1);

auto future = std::async(std::launch::async,
[&](){
//the previous allocation has been "freed" already, but is still in use
float* d_data2;
allocator.DeviceAllocate((void**)&d_data2, sizeof(float) * N, stream2);

evilzerokernel<<<1,1,0,stream2>>>(d_data2, 32);

allocator.DeviceFree(d_data2);

cudaStreamSynchronize(stream2);
}
);

future.wait();

cudaError_t status = cudaDeviceSynchronize();
assert(status == cudaSuccess);
}
```

With `cudaStreamPerThread`, the allocation is shared between both threads causing an error.

```
Device 0 allocated new device block at 0x7faebd800000 (4194304 bytes associated with stream 2).
0 available blocks cached (0 bytes), 1 live blocks outstanding(4194304 bytes).
Device 0 returned 4194304 bytes from associated stream 2.
1 available blocks cached (4194304 bytes), 0 live blocks outstanding. (0 bytes)
Device 0 reused cached block at 0x7faebd800000 (4194304 bytes) for stream 2 (previously associated with stream 2).
0 available blocks cached (0 bytes), 1 live blocks outstanding(4194304 bytes).
Device 0 returned 4194304 bytes from associated stream 2.
1 available blocks cached (4194304 bytes), 0 live blocks outstanding. (0 bytes)
error position 30, 0.000000 0.694820
```

If explicit streams are used instead, the program behaves as expected.

```
Device 0 allocated new device block at 0x7f6bbb800000 (4194304 bytes associated with stream 94428507066928).
0 available blocks cached (0 bytes), 1 live blocks outstanding(4194304 bytes).
Device 0 returned 4194304 bytes from associated stream 94428507066928.
1 available blocks cached (4194304 bytes), 0 live blocks outstanding. (0 bytes)
Device 0 allocated new device block at 0x7f6bbbc00000 (4194304 bytes associated with stream 94428507067088).
1 available blocks cached (4194304 bytes), 1 live blocks outstanding(4194304 bytes).
Device 0 returned 4194304 bytes from associated stream 94428507067088.
2 available blocks cached (8388608 bytes), 0 live blocks outstanding. (0 bytes)
```