NVIDIA / NVIDIA/cccl

[BUG]: Memory access error for thrust::sort

Open
#1,095 7 comments 1 reaction 1 assignee Claimed by @elstehle View on GitHub
needs triage
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Type of Bug

Runtime Error

### Component

Thrust

### Describe the bug

When compiling the following code https://gist.github.com/PhilipDeegan/ac3931cccb48e0e61e742fc2b731c4c7

on ubuntu 22.04

with command

```
clang++-17 --std=c++17 -fPIC -fsized-deallocation -Wno-unknown-cuda-version -Xclang -fcuda-allow-variadic-functions -x cuda --cuda-gpu-arch=sm_80 -O3 -o sorting_thrust sorting_thrust.cpp
```

I see

```
terminate called after throwing an instance of 'thrust::system::system_error'
what(): merge_sort: failed to synchronize: cudaErrorIllegalAddress: an illegal memory access was encountered
```

However:
I do not see this issue with thrust via AMD.
I do not see the issue with `-O0`
I do not see the issue if I compile with the define `-DDO_NOT_ADD_EB`
I continue to see no issues with the define `-DDO_NOT_ADD_EB` even with 2e7 items
And I do not see issues with only 1e3 items.

### How to Reproduce

Compile with previous instructions

### Expected behavior

No memory access error

----

godbolt doesn't support cuda clang 17`

### Reproduction link

_No response_

### Operating System

Ubuntu 22.04

### nvidia-smi output

Mon Nov 13 22:44:28 2023
+---------------------------------------------------------------------------------------+
| NVIDIA-SMI 545.23.06 Driver Version: 545.23.06 CUDA Version: 12.3 |
|-----------------------------------------+----------------------+----------------------+
| 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 A100-PCIE-40GB Off | 00000000:21:00.0 Off | 0 |
| N/A 69C P0 246W / 250W | 1171MiB / 40960MiB | 99% Default |
| | | Disabled |
+-----------------------------------------+----------------------+----------------------+

+---------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=======================================================================================|
+---------------------------------------------------------------------------------------+

### NVCC version

```
nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Fri_Sep__8_19:17:24_PDT_2023
Cuda compilation tools, release 12.3, V12.3.52
Build cuda_12.3.r12.3/compiler.33281558_0
```

```
clang++-17 -v
Ubuntu clang version 17.0.5 (++20231113073440+094cfd14c813-1~exp1~20231113073538.63)
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin
Found candidate GCC installation: /usr/bin/../lib/gcc/x86_64-linux-gnu/11
Selected GCC installation: /usr/bin/../lib/gcc/x86_64-linux-gnu/11
Candidate multilib: .;@m64
Selected multilib: .;@m64
Found CUDA installation: /usr/local/cuda-12.3, version
```

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.