RAII-style memory management

[1289941679]

Hello, I wonder whether this code is safe, and if it is, why?

void fn() {
    thrust::device_vector<int> d_data;
    
    // ... some data initialization

    kernel<<<...>>>(d_data.data(), d_data.size());
}

int main() {
    fn();
    // Now fn() returns, but is memory in d_data still safe? 
    // Will device memory in d_data freed and the still-running kernel read in invalid value?

    return 0;
}

[Curefab]

RAII-style means it is destructed at the end of the scope.

You write “now fn() returns”, but the comment is still in a line within the scope, before returning. Perhaps you can elaborate, if you mean after returning or at the position of the comment (sorry to be pedantic).

[1289941679]

Thanks! When function returns, d_data should be destroyed, device memory in d_data should be freed. However when the kernel is still running, will the device read in invalid value?

[Robert_Crovella]

presumably, thrust does something like cudaFree() when the device allocation goes out of scope.

regarding cudaFree() it seems evident that it does not (actually) free an allocation until device activity has completed.

To respond to a question posed in the comments:

@1289941679

// Now fn() returns, but is memory in d_data still safe?

When I read your question, I had the same questions that Curefab expressed. It seems like an odd code construct to begin with, and its not really clear at what point you are expressing a concern. Regarding that question, no: when fn() returns, memory in d_data is no longer “safe” (might be no longer safe/cannot be reliably counted out to continue to be safe).

In so far as the code execution within the curly-braces, I don’t see an issue.

If you were asking about the visibility of the device vector container at any point outside the curly-braces, just like std::vector, that would be inadvisable (out of scope).