using thrust device_vector as global variable

Question

Why does the following code crash at the end of the main?

#include <thrust/device_vector.h>

thrust::device_vector<float4> v;

int main(){
    v.resize(1000);
    return 0;
}

The error is:

terminate called after throwing an instance of 'thrust::system::system_error'
what():  unspecified driver error

If I use host_vector instead of device_vector the code run fine.

Do you think it's a Thrust bug, or am I doing something wrong here?

I tried it on ubuntu 10.10 with cuda 4.0 and on Windows 7 with cuda 6.5. The Thrust version is 1.7 in both cases.

thanks

Answer 1

The problem is neither a bug in Thrust, nor are you doing something wrong. Rather, this is a limitation of the design of the CUDA runtime API.

The underlying reason for the crash is that the destructor for the thrust::vector is being called when the variable falls out of scope, which is happening after the CUDA runtime API context has been torn down. This will produce a runtime error (probably cudaErrorCudartUnloading) because the process is attempting to call cudaFree after it has already disconnected from the CUDA driver.

I am unaware of a workaround other than not using Thrust device containers declared at main() translation unit scope.