Insufficient temporary memory allocated for device histograms

An internal user has reported a bug in `cub::DeviceHistogram`. When using 16-bit values, the computed temporary storage buffer size is too small on Pascal, leading to runtime errors. They've applied a workaround (shown below) that prevents the issue.

```
if (N == 1)
{
cub::DeviceHistogram::HistogramRange(
pTempStorage,
nTempStorageBytes,
reinterpret_cast(NULL),
pHistograms,
nLevels,
pLevels,
oROI.width,
oROI.height,
nRowStrideBytes,
hStream);

// Workaround for CUB 1.9.10 scratch buffer size bug on Pascal.
*hpBufferSize = static_cast(nTempStorageBytes + nTempStorageBytes + nLevels * sizeof(LevelType));
}
else if (N == 3)
{
cub::DeviceHistogram::MultiHistogramRange<3, 3>(
pTempStorage,
nTempStorageBytes,
reinterpret_cast(NULL),
pHistograms,
aLevels,
pLevels,
oROI.width,
oROI.height,
nRowStrideBytes,
hStream);

// Workaround for CUB 1.9.10 scratch buffer size bug on Pascal.
*hpBufferSize = static_cast(nTempStorageBytes + nTempStorageBytes + nLevels * sizeof(LevelType));
}
else
{
cub::DeviceHistogram::MultiHistogramRange<4, 4>(
pTempStorage,
nTempStorageBytes,
reinterpret_cast(NULL),
pHistograms,
aLevels,
pLevels,
oROI.width,
oROI.height,
nRowStrideBytes,
hStream);

// Workaround for CUB 1.9.10 scratch buffer size bug on Pascal.
*hpBufferSize = static_cast(nTempStorageBytes + nTempStorageBytes + nLevels * sizeof(LevelType));
}
}
```

More background in NVBug 200559924.