BlockReduce<int, ...> causes "an illegal instruction was encountered"
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
Consider the following kernel:
```cpp
template
__global__
void cudaMinMaxAmplitude(cudaTextureObject_t tex, int width, int height, int startRow, int *min, int *max) {
using ScalarT= uint8_t;
using BlockReduceT = BlockReduce;
__shared__ typename BlockReduceT::TempStorage minStorage;
__shared__ typename BlockReduceT::TempStorage maxStorage;
const auto col = blockIdx.x * blockDim.x + threadIdx.x;
const auto row = blockIdx.y * blockDim.y + threadIdx.y + startRow;
auto ampMin= ScalarT{};
auto ampMax= ScalarT{};
if (width > col && height > row) {
const auto [a, b, g, r] = tex2D(tex, col, row);
const auto amp = ScalarT(g);
ampMin = ampMax = amp;
} else {
ampMin = std::numeric_limits::max();
ampMax = std::numeric_limits::min();
}
auto blockMin = BlockReduceT(minStorage).Reduce(ampMin, cub::Min());
auto blockMax = BlockReduceT(maxStorage).Reduce(ampMax, cub::Max());
if (0 == threadIdx.x && 0 == threadIdx.y) {
atomicMin(min, blockMin);
atomicMax(max, blockMax);
}
}
```
If I change the first line to `using ScalarT= int;`, I see a "an illegal instruction was encountered" exception. Why?
Contributor guide
Assessment
This issue has not been assessed yet.