I have cuda code that uses thrust that I compile with compute_89,sm_89. Since upgrading my PC from an RTX 4090 to an RTX 5090 I now have these runtime errors:
CUDA error 209 [C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.9\include\cub/util_device.cuh, 101]: no kernel image is available for execution on the device
CUDA error 101 [C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.9\include\cub/util_device.cuh, 378]: invalid device ordinal
CUDA error 101 [C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.9\include\cub/device/dispatch/dispatch_for.cuh, 184]: invalid device ordinal
That seems to be caused by ‘thrust::for_each_n’. I tried to compile with 120 rather than 89 as the target also, and then it has a compile error about some data type sizes not matching.
I have the latest driver and the Cuda 12.9.1 SDK installed.
switching to compiling for compute_120,sm_120 is the correct approach to address the no kernel image error. (as is normal, it may also be possible to specify sm_89 along with an appropriate PTX definition/inclusion, such as -gencode arch=compute_89,code=compute_89 or -arch=sm_89)
If you have thrust code that compiles correctly for cc 8.9 but does not compile correctly for cc 12.0, you could file a thrust issue.
I tried compiling a simple example on godbolt using thrust::for_each_n on CUDA 12.8.1 for cc 12.0 and didn’t seem to have any compilation issues. (There is a runtime issue there, but its associated with the fact that the GPUs that godbolt uses are not cc 12. 0)
1>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.9\include\cuda/__ptx/instructions/generated/clusterlaunchcontrol.h(78): error : asm operand type size(4) does not match type/size implied by constraint ‘l’
1> : “l”((reinterpret_cast<long2>(&__try_cancel_response)).x),
1> ^
There isn’t much I can think of without a full test case to look at. Just to repeat myself, if you want to provide a full test case, you could file a thrust issue. You could also try the full test case I provided via godbolt, to see if it compiles correctly on your machine. That could be an instructive datapoint.
If your code compiles correctly on godbolt (i.e. cuda 12.8.1) then I would probably file a thrust issue. I doubt I will have any more observations.
I suppose a corrupted CUDA install might be another issue, so a clean reinstall might be worth a try, and you could also try updating your cccl to the latest. Bugs get fixed all the time.
I should note I tried reinstalling the Cuda 12.9.1 SDK (Which I assume put in the latest cccl) just now too (On Windows 10 64 Bit, Latest VS), and still the same error.
It seems like it could be a thrust issue if there was a change in thrust from 12.8.1 to 12.9.x
I’m not really sure what else it could be, at this point. Perhaps try compiling your code in godbolt. If it compiles there (which I would expect based on the tests you have run so far), then it is either a thrust issue or a machine setup issue.
You can file a bug also of course. But since thrust is part of an open source library, filing issues is the community-friendly way to report such things.
That was one of the reasons I thought the godbolt datapoint could be useful. Since the code I provided compiles there but not on your machine it points to either a toolchain issue (bug) or a machine setup issue.
It seems likely based on the data so far that CUDA 12.8.1 might unblock you on this issue.
A toolchain issue could be a change in the cccl libraries or it could be just a defect in the nvcc compiler toolchain, or perhaps some other issue. It would require more study to try to guess at whether it is a cccl library/header issue or nvcc toolchain bug. You could diff the relevant header files (where the asm error is reported) to see if there appear to be meaningful differences lately in the cccl headers.
But since you have found (or using my example) thrust code that seems to compile correctly on 12.8.1 and not on 12.9 (it seems/for whatever reason) then putting that in a thrust issue is one way to drive resolution and also makes the issue more visible to the community.
I wasn’t sure if there were any of the usual cmake custom build hoops to jump through at first, but
in the end I replaced these from the latest cccl github: