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.

Sign up for GitHub

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

Jump to bottom

[BUG]: Cub's ReduceByKey passing unexpected values to reduction_op #3890

Open
1 task done
sarda-devesh opened this issue Feb 21, 2025 · 1 comment
Open
1 task done

[BUG]: Cub's ReduceByKey passing unexpected values to reduction_op #3890

sarda-devesh opened this issue Feb 21, 2025 · 1 comment
Labels
bug Something isn't working right.

Comments

@sarda-devesh
Copy link

Is this a duplicate?

Type of Bug

Runtime Error

Component

CUB

Describe the bug

I was trying to utilize the cub::DeviceReduce::ReduceByKey function and noticed that my implementation was resulting in cudaGetLastError returning an illegal memory access error. Specifically, in my program, the values passed into the Reduce represent indexes, and the custom reduction function would use those passed in indexes to read values from a buffer. However, I noticed that for some threads in the last scheduled thread block, this function would pass in garbage values into this reduction function, causing it to access indexes greater than the buffer's size, resulting in illegal memory accesses.

How to Reproduce

You can find an example code recreating this behaviour at: https://godbolt.org/z/1eE96EYbo. When I compile and run this code using the command: nvcc cub_bug.cu -o cub_bug && ./cub_bug, I get the output that:

Running experiment with num records of 1048576
First 5 Key records: 0 1 2 3 4 
First 5 Value records: 0 0 0 0 0 
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [43,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [44,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [45,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [46,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [47,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [48,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [49,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [50,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [51,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [52,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [53,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [54,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [55,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [56,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [57,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [58,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [59,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [60,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [61,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [62,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [63,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [96,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [97,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [98,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [99,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [100,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [101,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [102,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [103,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [104,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [105,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [106,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [107,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [108,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [109,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [110,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [111,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [112,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [113,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [114,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [115,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [116,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [117,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [118,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [119,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [120,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [121,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [122,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [123,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [124,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [125,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [126,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [127,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [64,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [65,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [66,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [67,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [68,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [69,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [70,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [71,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [72,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [73,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [74,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [75,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [76,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [77,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [78,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [79,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [80,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [81,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [82,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [83,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [84,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [85,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [86,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [87,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [88,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [89,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [90,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [91,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [92,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [93,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [94,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [95,0,0] Assertion `lower_value < num_records` failed.
CUDA error: device-side assert triggered

Note that I am running this code on a AWS g6.xlarge instance with an Ubuntu 22.04 image and an L4 GPU. I have included the result of the nvidia-smi command with the cuda and driver versions.

Expected behavior

The CustomReductionOperator should never have an invalid index/value passed to it. For now, I have just added a check in my code to ignore any invalid values and I see the right result being produced but ideally, I shouldn't need to have a special case like this.

Reproduction link

https://godbolt.org/z/1eE96EYbo

Operating System

Ubuntu Linux 22.04

nvidia-smi output

 NVIDIA-SMI 570.86.10              Driver Version: 570.86.10      CUDA Version: 12.8     |
|-----------------------------------------+------------------------+----------------------+
| GPU  Name                 Persistence-M | Bus-Id          Disp.A | Volatile Uncorr. ECC |
| Fan  Temp   Perf          Pwr:Usage/Cap |           Memory-Usage | GPU-Util  Compute M. |
|                                         |                        |               MIG M. |
|=========================================+========================+======================|
|   0  NVIDIA L4                      Off |   00000000:31:00.0 Off |                    0 |
| N/A   29C    P8             15W /   72W |       1MiB /  23034MiB |      0%      Default |
|                                         |                        |                  N/A |
+-----------------------------------------+------------------------+----------------------+
                                                                                         
+-----------------------------------------------------------------------------------------+
| Processes:                                                                              |
|  GPU   GI   CI              PID   Type   Process name                        GPU Memory |
|        ID   ID                                                               Usage      |
|=========================================================================================|
|  No running processes found                                                             |

NVCC version

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Jan_15_19:20:09_PST_2025
Cuda compilation tools, release 12.8, V12.8.61
Build cuda_12.8.r12.8/compiler.35404655_0
@sarda-devesh sarda-devesh added the bug Something isn't working right. label Feb 21, 2025
@github-project-automation github-project-automation bot moved this to Todo in CCCL Feb 21, 2025
@lilohuang
Copy link

This has been a known issue for a while: #459

@sarda-devesh sarda-devesh mentioned this issue Feb 21, 2025
Open
6 tasks
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
bug Something isn't working right.
Projects
CCCL
Status: Todo
Milestone
No milestone
Development

No branches or pull requests
2 participants
@lilohuang @sarda-devesh