Shouldn't __threadfence() guarantee correct order of global writes?

I made a lock in cuda which protects the array "elements" from being a victim of a race condition. More precisely, only one block should be able to initialize elements[0].

//Bar.cu
Bar::Init(){
    while (true) {
        if (atomicExch(&lock, 1u) == 0u) { //<-- my lock
           
           //only one block should initialize
           if(elements[0] == 0){

                //create a new element
                Foo temp = Foo();
                volatile Foo* new_element = malloc(sizeof(Foo));
                memcpy((void*)new_element, &temp, sizeof(Foo));

                //global write
                elements[0] = new_element; 
           }
         
           
           __threadfence(); //<-- make sure next block sees change
           atomicExch(&lock, 0u);
           break;
        }
    }
}

I know __threadfence() is suppose to make the order of global writes guaranteed as postulated by CygnusX1. Therefore, I presume elements[0] = new_element is written to global memory before atomicExch(&this->lock, 0u), such that the next block which aquires the lock will now see elements[0] as new_element, not as 0.

However, when I run the code, the next block doesn't always register that elements[0] has been updated. Why? Shouldn't the fact that I defined elements with "volatile" make all reads of elements be fetched from global memory, not cache? Then why doesn't __threadfence() do the job here?

The rest of my code:

//Bar.cuh
class Bar{
   public:
       volatile Foo ** elements;
       unsigned int lock = 0u;
       __device__ void Init();

}

\\Main.cu

__global__ void MyKernel(Bar* bar){
    bar->Init();
}

int main(int argc, char* argv[])
{
    Bar bar = Bar();
    cudaMalloc((void**)&bar.elements, 10 * sizeof(Foo)); //<--room for 10 elements
    cudaMemset(bar->elements, 0, 10 * sizeof(Foo*)); //<-- set all elements to 0
    Bar * bar_pointer = new Bar();
    cudaMalloc((void**)&bar_pointer , sizeof(Bar));
    cudaMemcpy(bar_pointer , &bar, sizeof(Bar),cudaMemcpyHostToDevice);
    MyKernel<< <2,1>> > (bar_pointer); // <-- launch kernel with two blocks
}