Can I 'swap' memory between __shared__ and __global__ (or __global__ and __global__) arrays?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I’m curious if some of the problems I’m seeing in a branch of my code might stem from the fact that I am trying to “swap” memory from arrays in __shared__ to arrays in __global__ memory. The idea is to have the thread block loop over each index of the respective arrays and do the typical swap operation:

__shared__ float sh_array[N_ELEMENTS];
float* gbl_array;

int pos = threadIdx.x;
while (pos < N_ELEMENTS) {
  float tmp = sh_array[pos];
  sh_array[pos] = gbl_array[pos];
  gbl_array[pos] = tmp;
  pos += blockDim.x;
}

Will that require synchronization to get it right? It seems that I may need to do the following:

__shared__ float sh_array[N_ELEMENTS];
float* gbl_array;

int base_pos = 0;
while (base_pos < N_ELEMENTS) {
  const int pos = base_pos + threadIdx.x;
  float tmp;
  if (pos < N_ELEMENTS) {
    tmp = sh_array[pos];
    sh_array[pos] = gbl_array[pos];
  }
  __syncthreads();
  if (pos < N_ELEMENTS) {
    gbl_array[pos] = tmp;
  }
  __syncthreads();
  base_pos += blockDim.x;
}

Can anyone confirm that the problem I think I’m seeing is real and the solution looks sound?

2

Update: I have implemented this solution for my main code, and it seems to fix the problem. I’m not sure if things are completely fine, but this definitely cleaned up some otherwise inexplicable pollution in my numbers. I did not end up needing the second __syncthreads() in the pseudo-code above, as the threads will move on to new indices of the arrays in subsequent iterations of the loop. In my actual code, I do have a __syncthreads() further down to catch the tail end of the writes to ensure that the swap is complete once it’s time to swap back.

Thank you, CUDA engineers! __syncthreads() and __syncwarp() memory barriers are truly impressive.

3

If each block operates on a separate gbl_array, you don’t need any synchronization because different threads access distinct array positions.
If multiple blocks have the same gbl_array, you have a race condition between blocks which cannot be fixed by __syncthreads()

1 Like
4

What if I have different blocks operating on the same global array, but different sectors of it? Like, block 0 operates on elemetns 0 through 1023, block 1 operates on elements 1024 through 2048… I think that’s OK. In effect, the gbl_array pointer is unique to each block, although if a block overruns its bounds (which I am careful to prevent) then yes, __syncthreads() would not help me with that, I agree.

But I did seem to fix a lot of problems when I began adding the synchronization between reads from some location in global memory and writes to it, even though any given address is only operated upon by a single thread. Are you sure tha no synchronization should be needed?

5

Yes when all blocks are working with their own non-overlapping sections its fine. (It’s not obvious from the code snippet since all blocks would use the same pos when only threadIdx.x is used.

I am sure that if an address is only accessed by a single thread there cannot be a data race for this address. No synchronization required.

1 Like
6

Yes, sorry I didn’t make the distinct sectors clear in my example. If this is true, then, I need to review what I’ve done, because I may have merely created a Heisenbug. It certainly went from “very bad, always bad” to “pretty quiet” as soon as I did what I describe above, and the new synchronization isn’t really much most costly, so far as I can tell, than the old one, so I don’t think I’ve changed the code in a way that would drastically tamp down on some other collision that’s happening. I will continue to investigate…

7

Have you run this code with compute-sanitizer? It can find many instances of race conditions, though not all of them. It also seems possible that you could have an out-of-bounds access somewhere which may not be obvious if it is an off-by-one error that does not trigger a memory access violation. Adding __sync_threads() may be merely masking the root cause of the observed data corruption in most cases, as you already suspected yourself.

1 Like
8

I have been running compute-sanitzer, but I will do more of that.

9

@striker159 @njuffa Thanks for insisting. Indeed, I just removed the excess synchronization, and the code continues to produce identical results after a 21-minute run. All that really happened is that the code finishes about 5% faster.

10

CLosing this topic as I found the (very sneaky) race condition that was getting in the way. Marching on!

11

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.