Is &p->member well-defined for a device pointer p in CUDA host code?

Suppose I have the following struct, with an array followed by a counter, where the size of the array is potentially big: struct myStruct { int values[N]; int counter; } and a producer-consumer setup. I see three options for resetting the counter (take the address of the counter member, calculate that address, or reset the entire structure). myStruct *device_ptr; cudaMalloc(&device_ptr, sizeof(*device_ptr)); myStruct *host_ptr; cudaHostAlloc(&host_ptr, sizeof(*host_ptr), 0); while(condition) { /* reset counter */ // Option A: cudaMemset(&device_ptr->counter, 0, sizeof(device_ptr->counter)); // Option B: cudaMemset(reinterpret_cast(device_ptr) + offsetof(myStruct, counter), 0, sizeof(device_ptr->counter)); // Option C: cudaMemset(device_ptr, 0, sizeof(*device_ptr)); myKernel<<>>(device_ptr); cudaMemcpy(host_ptr, device_ptr, sizeof(*device_ptr), cudaMemcpyDeviceToHost); process(host_ptr); } Option C should be the most straight-forward, but for big N maybe not ideal if the producer only inserts values based on the counter. My understanding is that option B is also well-defined, since it boils down to compile-time pointer arithmetic. However from a readability (at least for me) standpoint, I'd prefer option A. It would be well-defined in C++, but I couldn't find anything specific regarding the CUDA context I am in. Although NVCC appears to have no issue in a small test program, looking at the programming guide, there is this, which lets me to believe it is undefined behaviour and at most compiler dependent: Pointer dereferencing (*pointer, pointer->member, pointer[0]) is allowed only in the same execution space where the associated memory resides.