# CUDA 8 - Thrust bug(?) **URL:** \ **Category:** CUDA Programming and Performance\ **Created:** [September 29, 2016, 3:02pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235 "2016-09-29T15:02:16Z")\ **Posts on this page:** 17\ **Page:** 1
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [September 29, 2016, 3:02pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/1 "2016-09-29T15:02:16Z")
Since CUDA 8, I am having trouble with Thrust. I am using “sort” kernel from driver API, to sort a buffer of data according to some keys. The whole kernel is as follows: ```auto #include #include #include "CudaCommon.h" //some defines, basic math, etc... MyDataType is defined there extern "C" __global__ void sortKernel( uint64_t* keys, MyDataType* data, unsigned int dataSize ) { thrust::sort_by_key(thrust::device, keys, keys + dataSize, data); } ``` size of MyDataType is 48 bytes (12 ints). The kernel itself is launched just as one instance (groups & block size equaling to 1). ```auto checkCudaErrors(cuLaunchKernel(_sortKernel, 1, 1, 1, 1, 1, 1, 0, 0, sortArgs, nullptr)); checkCudaErrors(cuEventRecord(_kernelSyncEvent, 0)); checkCudaErrors(cuEventSynchronize(_kernelSyncEvent)); ``` This code works OK on CUDA 7.5, on CUDA 8 (RC and Release) it causes CUDA\_ERROR\_UNKNOWN (on the cuEventSynchronize). System specs: W10 x64, i7 4770K, VS2015 (without any further update), 16GB RAM, GTX 780, drivers 369.30 (shipped with CUDA 8), CUDA installer “cuda\_8.0.44\_win10.exe”. What is wrong? An error in my code, or a bug in Thrust? Thanks in advance. ---
**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\ **Post date:** [September 29, 2016, 4:52pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/2 "2016-09-29T16:52:50Z")
If you want to provide a complete app that I can compile and run and see the error, I’ll take a look as time permits. What is the purpose of calling sort from the kernel as you have? Why not call it directly from host code? ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [September 29, 2016, 5:05pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/3 "2016-09-29T17:05:36Z")
Sorry, cannot provide the whole app. I think the data I’ve provided is pretty much sufficient, the kernel in dispute is running just the thrust code, nothing else. And the data to be sorted is literally just 12 ints wrapped together. BTW I have not found a way how to call Thrust::sort from driver API directly, that’s why it’s wrapped in a kernel… ---
**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\ **Post date:** [September 29, 2016, 5:26pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/4 "2016-09-29T17:26:32Z")
Thrust depends on the runtime API. You can intermix driver API and runtime API into the same project. You cannot use thrust from host code in a driver API-only project. It requires linking with the runtime API libraries, at a minimum. ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [September 29, 2016, 5:36pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/5 "2016-09-29T17:36:05Z")
Well, that’s why I am calling it from a kernel, that seems to work OK… at least did until now… ---
**Author:** ![MutantJohn](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MutantJohn](https://forums.developer.nvidia.com/u/MutantJohn)\ **Post date:** [September 29, 2016, 10:15pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/6 "2016-09-29T22:15:25Z")
> [@](#): > > You cannot use thrust from host code in a driver API-only project. It requires linking with the runtime API libraries, at a minimum. Yeah, how are you compiling your code, Jofo? I had weird and seemingly random Thrust linking errors because I wasn’t linking against cudadevrt. ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [September 30, 2016, 6:03am UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/7 "2016-09-30T06:03:50Z")
In CMAKE (debug info turned on…) ```auto CUDA_COMPILE_PTX(PTX_FILES ${SRC_CUDA} OPTIONS -arch compute_35 -lineinfo -G) ``` SRC\_CUDA contains all the kernels I have in my project. This is then followed by a bunch of scripts copying the kernels where I need em, wrapped in a custom target. ---
**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\ **Post date:** [September 30, 2016, 1:18pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/8 "2016-09-30T13:18:03Z")
agree with MutantJohn, calling thrust from device code may necessitate additional compile requirements. You can look at any CUDA sample project that calls kernels from device code to get an idea of the needed switches. ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [September 30, 2016, 11:28pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/9 "2016-09-30T23:28:05Z")
This code was working well in CUDA 7.5, but no longer does in CUDA 8… so my linking must be fine (I guess) ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [October 1, 2016, 1:21pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/10 "2016-10-01T13:21:41Z")
OK I’ve run into a weird situation… I’ve hacked together a simple demo project which uses similar mechanics as my production code, but now it throws CU\_ERROR\_ILLEGAL\_INSTRUCTION even on CUDA 7.5, whereas my production code works… although in debug, the sorting kernel runs really slow even on let’s say 35000 items. > **[MEGA](https://mega.nz/#F!QhQSkbTS!vEHRk3S_OdzEQNHD5EITVA)** > > MEGA provides free cloud storage with convenient and powerful always-on privacy. Claim your free 20GB now Screw it, I will probably switch to a different library for sorting, any ideas? ---
**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\ **Post date:** [October 1, 2016, 1:58pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/11 "2016-10-01T13:58:27Z")
Side remark: For CUDA debug builds, the compiler turns off all optimizations, and the resulting code may run ten times slower than the code from a release build, which defaults to -O3, that is, maximum optimization. ---
**Author:** ![MutantJohn](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MutantJohn](https://forums.developer.nvidia.com/u/MutantJohn)\ **Post date:** [October 1, 2016, 4:00pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/12 "2016-10-01T16:00:12Z")
I wouldn’t just abandon thrust. Looking at what you have now, that’s not a strong candidate for dynamic parallelism. Just call it from the host. ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [October 1, 2016, 4:22pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/13 "2016-10-01T16:22:18Z")
I cannot call it from the host because I use driver API… ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [October 1, 2016, 4:31pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/14 "2016-10-01T16:31:01Z")
Yeah, I forgot about the debug flag, that sorts the speed… Moreover - correct me if I am wrong, but I think Thrust manages the kernel execution (“dyamic parallelism”) by itself in terms of grid/block size, so it should be sufficient to just call 1 thread which calls the whole subkernel. Or not? ---
**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\ **Post date:** [October 1, 2016, 4:50pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/15 "2016-10-01T16:50:16Z")
> [@](#): > > Moreover - correct me if I am wrong, but I think Thrust manages the kernel execution (“dyamic parallelism”) by itself in terms of grid/block size, so it should be sufficient to just call 1 thread which calls the whole subkernel. Or not? When you call thrust algorithms from device code using thrust::device execution policy, you are leaving the decision up to thrust as to how best to run this code. It means that: 1. the implementation thrust chooses may vary depending on target architecture and compilation switches 2. the thrust implementation may vary from thrust version to thrust version (note that you can override this ambiguity using other thrust execution policies) ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [October 2, 2016, 11:16am UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/16 "2016-10-02T11:16:00Z")
A little more update I’ve put this as a bug on nV support, and the reponse so far - first of all, if I want to use dynamic parallelism, I need to compile the ptx with -rdc=true flag, otherwise it would run sequentially (I supposed it means in a single GPU thread?). But doing this requires linking against cudadevrt.lib using cuLink, which would be insanely stupid - I would then need a path to the library in run time. An example is actually in the documentation here: [http://docs.nvidia.com/cuda/nvrtc/index.html#example-dynamic-parallelism](http://docs.nvidia.com/cuda/nvrtc/index.html#example-dynamic-parallelism) So I wonder, if there is a possible solution using CMAKE and doing it in program compile time, I will try to experiment with it… - secondly, it is possible to mix runtime and driver API [http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#interoperability-between-runtime-and-driver-apis](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#interoperability-between-runtime-and-driver-apis) but from my initial experiments I’ve found out that this will not work with thrust, because when I modified the code from the example I had provided, it will not compile. I’ve included thrust and cuda\_runtime\_api.h into host, then called ```auto thrust::sort(thrust::device, (uint64_t*)g_bufferKeys, ((uint64_t*)g_bufferKeys) + numElements, (MyDataType*)g_bufferData); ``` nVidia guys even left a note in thrust/system/cuda/detail/sort.inl(203) that the code has to be NVCC’ed first. So this is probably a no-go. ---
**Author:** ![Jofo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jofo](https://forums.developer.nvidia.com/u/Jofo)\ **Post date:** [October 2, 2016, 12:58pm UTC](https://forums.developer.nvidia.com/t/cuda-8-thrust-bug/45235/17 "2016-10-02T12:58:30Z")
So I’ve altered the example case and put the -rtc=true flag to the PTX compilation, then linked with cudadevrt.lib in run time. It actually linked, but the kernel won’t run, the error being CUDA\_ERROR\_LAUNCH\_FAILED and there was a message in the console “temporary\_buffer::allocate: get\_temporary\_buffer failed”. Now what?