Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[BUG]: reduce_by_key fails with zip_iterator to const pointers #1527

Open
1 task done
upsj opened this issue Mar 9, 2024 · 5 comments
Open
1 task done

[BUG]: reduce_by_key fails with zip_iterator to const pointers #1527

upsj opened this issue Mar 9, 2024 · 5 comments
Labels
bug Something isn't working right.

Comments

@upsj
Copy link

upsj commented Mar 9, 2024

Is this a duplicate?

Type of Bug

Compile-time Error

Component

Thrust

Describe the bug

The following code fails to compile with CUDA 12.4 (tracked by upstream issue ginkgo-project/ginkgo#1564)

#include <thrust/execution_policy.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/reduce.h>
#include <thrust/tuple.h>

void sum_duplicates(int size) {
  const int *rows{};
  const int *cols{};
  const float *in_vals{};
  int *out_rows{};
  int *out_cols{};
  float *out_vals{};
  auto in_locs = thrust::make_zip_iterator(thrust::make_tuple(rows, cols));
  auto out_locs =
      thrust::make_zip_iterator(thrust::make_tuple(out_rows, out_cols));
  thrust::reduce_by_key(thrust::device, in_locs, in_locs + size, in_vals,
                        out_locs, out_vals);
}

Noteworthy: The issue goes away when the pointers are int* instead of const int*

How to Reproduce

  • call nvcc file.cu

Expected behavior

The code compiles (as it did with CUDA <= 12.3)

Reproduction link

No response

Operating System

Ubuntu 22.04

nvidia-smi output

No response

NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Tue_Feb_27_16:19:38_PST_2024
Cuda compilation tools, release 12.4, V12.4.99
Build cuda_12.4.r12.4/compiler.33961263_0

@upsj upsj added the bug Something isn't working right. label Mar 9, 2024
@upsj
Copy link
Author

upsj commented Mar 9, 2024

A possible workaround here is to const_cast the inputs, but I would like to hear if there is a less "unsafe" workaround we can apply.

@miscco
Copy link
Collaborator

miscco commented Mar 10, 2024

Are you building against current main or a specific branch?

@upsj
Copy link
Author

upsj commented Mar 10, 2024

I just checked, the issue is still present in main

miscco added a commit to miscco/cccl that referenced this issue Mar 11, 2024
It seems that explicitly passing in `thrust::device` was key here, otherwise the bug did not manifest

Fixes [BUG]: `reduce_by_key` fails with zip_iterator to const pointers NVIDIA#1527
@miscco
Copy link
Collaborator

miscco commented Mar 11, 2024

@upsj Thanks a lot for the bug report. I opened a PR with a fix and added your reproducer to out test suite

miscco added a commit to miscco/cccl that referenced this issue Mar 11, 2024
It seems that explicitly passing in `thrust::device` was key here, otherwise the bug did not manifest

Fixes [BUG]: `reduce_by_key` fails with zip_iterator to const pointers NVIDIA#1527
@upsj
Copy link
Author

upsj commented Mar 11, 2024

That was quick, thank you very much!

wmaxey pushed a commit that referenced this issue Mar 12, 2024
* Ensure that we can run `reduce_by_key` with const inputs

It seems that explicitly passing in `thrust::device` was key here, otherwise the bug did not manifest

Fixes [BUG]: `reduce_by_key` fails with zip_iterator to const pointers #1527

* Also address nvbug4550097

The explicit usage of `int` can give a conversion warning, so just use the right difference type

* Update thrust/testing/zip_iterator_reduce_by_key.cu
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working right.
Projects
Status: Todo
Development

No branches or pull requests

2 participants