# terminate called after throwing an instance of 'thrust::system::system_error' what():  after cudaGetLastError: an illegal memory access was encountered

- **ID:** `cuda/thrust-vector-alloc-failure`
- **Domain:** cuda
- **Category:** runtime_error
- **Error Code:** `cudaErrorIllegalAddress (77)`
- **Verification:** ai_generated
- **Fix Rate:** 80%

## Root Cause

Thrust algorithms (e.g., thrust::sort, thrust::reduce) internally launch CUDA kernels; an illegal memory access occurs when a device vector is accessed out of bounds or a host pointer is passed to a device-side operation, causing a kernel crash that Thrust reports as a system_error.

## Version Compatibility

| Version | Status | Introduced | Deprecated |
|---------|--------|------------|------------|
| CUDA 11.8 | active | — | — |
| CUDA 12.1 | active | — | — |
| Thrust 1.17.2 | active | — | — |
| Thrust 2.1.0 | active | — | — |

## Workarounds

1. **Use cuda-memcheck or compute-sanitizer to pinpoint the exact illegal memory access location. Example: run the program with compute-sanitizer --tool memcheck ./my_program, then fix the out-of-bounds index.** (90% success)
   ```
   Use cuda-memcheck or compute-sanitizer to pinpoint the exact illegal memory access location. Example: run the program with compute-sanitizer --tool memcheck ./my_program, then fix the out-of-bounds index.
   ```
2. **Ensure all device vectors are properly sized and that host pointers are not accidentally passed to Thrust algorithms. Use thrust::device_vector for device data and thrust::host_vector for host data. Example: change thrust::sort(host_ptr, host_ptr+N) to thrust::sort(d_vec.begin(), d_vec.end()).** (85% success)
   ```
   Ensure all device vectors are properly sized and that host pointers are not accidentally passed to Thrust algorithms. Use thrust::device_vector for device data and thrust::host_vector for host data. Example: change thrust::sort(host_ptr, host_ptr+N) to thrust::sort(d_vec.begin(), d_vec.end()).
   ```

## Dead Ends

- **** — Resetting the device (cudaDeviceReset) does not fix the root cause; the illegal access will recur on the next Thrust call. (90% fail)
- **** — Increasing the device vector size arbitrarily may mask the out-of-bounds access but does not guarantee correctness; the real bug is in the algorithm logic. (75% fail)