Error when using thrust::fill on statically-allocated device array

Accelerated Computing CUDA CUDA Programming and Performance
1 
__device__ int A[4];

int main() {
    thrust::fill(
        thrust::device_ptr<int>(A),
        thrust::device_ptr<int>(A) + 4,
        1
    );
    return 0;
}

This simple example fails with the error:

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  trivial_device_copy D->H failed: an illegal memory access was encountered

This works:

int main() {
    int *A2;
    cudaMalloc((void **)&A2, 4 * sizeof(int));

    thrust::fill(
        thrust::device_ptr<int>(A2),
        thrust::device_ptr<int>(A2) + 4,
        1
    );
    return 0;
}
2

Ok, I’ve figured out I can use cudaGetSymbolAddress like so:

int *A_device;
cudaGetSymbolAddress((void **)&A_device, A);
thrust::fill(
    thrust::device_ptr<int>(A_device),
    thrust::device_ptr<int>(A_device) + 4,
    1
);

Is this the clearest way to express this in Thrust?

3

That looks to be the best implementation.

https://stackoverflow.com/questions/34932416/sorting-statically-allocated-array-using-thrust