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

@RobertCrovella: That isn't what is happening here. This is happening because the vector drops out of scope and is destroyed after the CUDA context is torn down. So effectively cudaFree is getting called with no working runtime API connection, resulting in a runtime error — talonmies

talonmies talonmies · Accepted Answer · 2014-09-01T08:03:22

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.

using thrust device_vector as global variable

1 Answers