Weird behavior of Thrust zip iterator
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
Hello,
I am seeing that thrust zip iterator is not working as expected. See the attached rerproducer.
```
#include
#include
#include
#include
#include
#include
#include
#define CUDA_TRY(call) \
do { \
cudaError_t ret = (call); \
if (ret != cudaSuccess) { \
std::cout << "CUDA error" << std::endl; \
std::abort(); \
} \
} while (0);
struct alignas(16) pair_t {
uint64_t a_;
float b_;
};
struct triplet_t {
std::size_t a_{};
std::size_t b_{};
std::size_t c_{};
};
struct dummy_member_a_t {
pair_t a_{};
int b_{};
triplet_t c_{};
int32_t* d_{};
};
struct dummy_member_b_t {
uint64_t a_{};
thrust::equal_to f_{};
};
struct dummy_member_t {
dummy_member_a_t a_{};
dummy_member_b_t b_{};
float c_{};
};
struct dummy_unused_t {
dummy_member_t a_{};
float b_{};
};
struct transform_key_t {
dummy_unused_t dummy_unused_{};
__device__ float operator()(uint64_t key) const { return static_cast(key * 5); }
};
template
__global__ void print_kernel(InputIterator input_first, InputIterator input_last)
{
auto const tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < thrust::distance(input_first, input_last)) {
printf("kernel print tid=%d key=%llu transformed=%f\n",
tid,
thrust::get<0>(*(input_first + tid)) /* this works OK */,
thrust::get<1>(*(input_first + tid)) /* this does not work OK */);
}
}
int main()
{
cudaStream_t stream{nullptr};
std::vector h_keys = {0, 1, 2, 3, 4};
uint64_t* d_keys{};
CUDA_TRY(cudaMallocAsync(&d_keys, h_keys.size() * sizeof(uint64_t), stream));
CUDA_TRY(cudaMemcpyAsync(
d_keys, h_keys.data(), h_keys.size() * sizeof(uint64_t), cudaMemcpyHostToDevice, stream));
CUDA_TRY(cudaStreamSynchronize(stream));
auto dummy_unused = dummy_unused_t();
auto transformed_first =
thrust::make_transform_iterator(d_keys, transform_key_t{dummy_unused});
auto input_first = thrust::make_zip_iterator(d_keys, transformed_first);
print_kernel<<<1, 256, 0, stream>>>(input_first, input_first + h_keys.size());
CUDA_TRY(cudaFreeAsync(d_keys, stream));
CUDA_TRY(cudaStreamSynchronize(stream));
return 0;
}
```
If I run this, I get
```
kernel print tid=0 key=0 transformed=0.000000
kernel print tid=1 key=1 transformed=10.000000 <= should be 5
kernel print tid=2 key=2 transformed=20.000000 <= should be 10
kernel print tid=3 key=3 transformed=0.000000 <= should be 15
kernel print tid=4 key=4 transformed=0.000000 <= should be 20
```
I compiled this with CUDA 11.8 (`/usr/local/cuda-11.8/bin/nvcc -arch=sm_86 --expt-relaxed-constexpr ./reproducer.cu`) and tested on RTX A6000.
This happens with CUDA 11.8, 11.7, and 12.2 but this does not happen with CUDA 11.6 (v11.6.124). And this weird behavior also disappears if I compiled the code with `-G`.
Contributor guide
Assessment
This issue has not been assessed yet.