NVIDIA / NVIDIA/cccl

BlockReduce<int, ...> causes "an illegal instruction was encountered"

Open
#934 3 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.