A safe CUDA __shared__ array container