NVIDIA / NVIDIA/cccl

[BUG]: `thrust::device_vector` does not compile with large element type

Open
#2,777 0 comments 0 reactions 0 assignees View on GitHub
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

Compile-time Error

### Component

Thrust

### Describe the bug

`thrust::device_vector` cannot be used with large element type (sizeof(T) ~ 32kb).
Any fill kernel or initialization kernel will fail to compile because constant memory space for kernel parameters are exceeded.

>/opt/compiler-explorer/cuda/12.6.1/bin/../targets/x86_64-linux/include/cub/device/dispatch/dispatch_for.cuh(137): Error: Formal parameter space overflowed (40016 bytes required, max 32764 bytes allowed) in function _ZN3cub17CUB_200500_890_NS6detail8for_each13static_kernelINS2_12policy_hub_t12policy_350_tEmN6thrust20THRUST_200500_890_NS8cuda_cub20__uninitialized_fill7functorINS7_10device_ptrI11LargeStructEESC_EEEEvT0_T1_

### How to Reproduce

```
#include

struct LargeStruct{
char data[40000];
};

int main(){
thrust::device_vector d_vec(1);
}
```

### Expected behavior

device_vector can be used with large types.

### Reproduction link

https://godbolt.org/z/88v1ffcYh

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_

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.