[warp's asynchrous behavior] does warp 0 always gets executed first?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, I was trying to study the asynchrous behavior of warp scheduling, and have created the following test code:

from numba import cuda
import torch

N = 128
a = torch.zeros(N, device='cuda', dtype=torch.float32)

@cuda.jit
def kernel(a, b, N):
    tx = cuda.threadIdx.x
    if tx == 0:
        a[0] = 100

    b[tx] = a[tx] + 1

for i in range(1000):
    b = torch.empty_like(a)
    kernel[1, N](a, b, N)
    if b[0] != 101:
        # Found a case where a[0] is not written to yet
        print(b)

Basically inside my kernel, I make thread 0 writes to a[0], but I expect other warps could be executing line b[tx] = a[tx] + 1, before thread 0 writes to a[0], if other warps get scheduled first.

However, even after running this code 1000 times, I didn’t observe one time that b[0] becomes 0 + 1, instead of 100 + 1, does this mean that when thread 0 writes to a[0], a[0] = 100 is guaranteed to happen before b[tx] = a[tx] + 1? I did try to make thread 64 writes to a[0] instead, and was able to observe b[tx] becomes 1 instead of 101.

Thanks!

2

The order of execution among warps is not defined by CUDA. Even if you observe the same behavior 1000 times.

More generally, the CUDA programming model does not define the order of execution of threads. If the programmer needs ordering for correctness, the programmer must provide for that via explicit methods such as synchronization, cooperative groups, etc.