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

[fkallen]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

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 <thrust/device_vector.h>

struct LargeStruct{
    char data[40000];
};

int main(){
    thrust::device_vector<LargeStruct> 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