cudaLaunchCooperativeKernel behaviour

Accelerated Computing CUDA CUDA Programming and Performance
1

I have a kernel that has groups of N threadblocks cooperate to iteratelively produce a result. They communicate via atomic counters in global memory (acquire/release pattern, each block of work has its own set of atomic counters). This works fine when I have exclusive access to the GPU and can expect at least N threadblocks to be resident at once.

However, it is entirely possible that other processes running simultaneously will occupy some GPU resources for a long time, this is out of my control. Other processes could even be trying to run a similar kernel. I worry that two or more processes could end up deadlocked if they launch the kernel with <N threadblocks resident.

Thus I want to use cudaLaunchCooperativeKernel to ensure a whole group of N threadblocks is resident at once.

In experiments I found that it’s possible to launch cooperative kernels on multiple streams and have them run concurrently. But behaviour is different on Hopper and Blackwell GPUs compared to Ada and Ampere. On the newer GPUs, the kernel will in fact only launch once there are enough resources available to have the entire grid resident at once. But on Ada/Ampere (tested on A100, L40S, RTX5000 Ada, RTX A6000) the driver will in fact launch partial grids, i.e. some threadblocks will start running before there is space for the entire grid to run on the GPU concurrently.

Is this expected behaviour? And how can I ensure that two processes won’t deadlock while waiting for GPU resources to become available?

2

That wouldn’t line up with my expectation, unless you are talking about some incredibly short interval, on the order of the block scheduling latency. Perhaps more description about how you reached that conclusion, or a complete example might be useful.

since you’ve advanced a statement as true that I don’t agree with, it’s not possible for me to comment further. If someone says to me “explain to me why the world is flat” I cannot offer any commentary there.

I will point out that the GPU handles this in the general case (not MIG, not MPS, not Exclusive-Process compute mode, not pre-pascal GPU) via preemption, as far as I know, otherwise described as context-switching. Context-switching or preemption could allow for the case where Process A launches work that utilizes most of the GPU, and yet a cooperative kernel in process B (requiring, lets say, the whole GPU) can begin running before the work launched by Process A has completed.

3

Here’s a minimal example that shows the behaviour. This launches four cooperative grids on four parallel streams. Each grid has (#SM / 3 + 2) threadblocks. Each threadblock occupies a full SM due to claiming the maximum amount of dynamic shared memory. So only two full grids will fit on the GPU at once.
The test kernel simply records, per threadblock, the time (%globaltimer) at launch, after a first grid sync, after waiting for a given number of milliseconds after launch, and after a second grid sync.

// cooperative_launch.cu
// Build and run with:
// $ nvcc -o cooperative_launch cooperative_launch.cu -arch=all
// $ CUDA_VISIBLE_DEVICES=0 ./cooperative_launch
#include <cooperative_groups.h>
#include <cuda.h>

#include <cstdint>
#include <cstdio>

struct Result {
    uint64_t t0;
    uint32_t t1, t2, t3;
    uint32_t sm, tb;
    int grid;
};

__device__ uint64_t globaltimer() {
    uint64_t t;
    asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(t)::"memory");
    return t;
}

__global__ void __launch_bounds__(1) test_kernel(Result *out, int grid, int offset, int ms) {
    uint64_t t0 = globaltimer();
    cooperative_groups::this_grid().sync();
    uint32_t t1 = globaltimer() - t0;
    uint32_t t2 = t1;
    while (t2 < ms * 1000000) {
        t2 = globaltimer() - t0;
    }
    cooperative_groups::this_grid().sync();
    uint32_t t3 = globaltimer() - t0;
    uint32_t smid;
    asm volatile("mov.u32 %0, %%smid;" : "=r"(smid));
    out[offset + blockIdx.x] = Result{t0, t1, t2, t3, smid, blockIdx.x, grid};
}

__global__ void print_results(Result *in, int nresults) {
    if (threadIdx.x + blockIdx.x == 0) {
        uint64_t t_min = in[0].t0;
        for (int i = 1; i < nresults; ++i) {
            t_min = min(t_min, in[i].t0);
        }
        for (int i = 0; i < nresults; ++i) {
            auto &r = in[i];
            printf("%3d Grid %2d, TB %2d, SM %3d - Launched at %lluns, grid sync after %uns, done "
                   "waiting after %uns, grid sync after %uns\n",
                   i, r.grid, r.tb, r.sm, r.t0 - t_min, r.t1, r.t2, r.t3);
        }
    }
}

int main() {
    constexpr int ngrids = 4;
    cudaStream_t stream[ngrids];
    for (int i = 0; i < ngrids; ++i) {
        cudaStreamCreateWithFlags(&stream[i], cudaStreamNonBlocking);
    }

    int device = 0, can_coop = 0, max_smem = 0, sm_count = 0;
    cudaGetDevice(&device);
    cudaDeviceGetAttribute(&can_coop, cudaDevAttrCooperativeLaunch, device);
    cudaDeviceGetAttribute(&max_smem, cudaDevAttrMaxSharedMemoryPerBlockOptin, device);
    cudaDeviceGetAttribute(&sm_count, cudaDevAttrMultiProcessorCount, device);
    // Use all smem to ensure a single TB per SM
    cudaFuncSetAttribute(test_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, max_smem);

    // Choose grid size so only two full grids will fit on the GPU, with some SMs idle.
    int tbs = (sm_count / 3) + 2;
    printf("SM count %d, can coop %d, max smem %d. Grids %d x %d\n", sm_count, can_coop, max_smem,
           ngrids, tbs);

    Result *bfr;
    cudaMalloc(&bfr, 4096 * sizeof(Result));
    for (int i = 0; i < ngrids; ++i) {
        int offset = i * tbs;
        int ms = (i == 0) ? 9 : 1;  // First grid runs for longer
        void *args[] = {&bfr, &i, &offset, &ms};
        cudaLaunchCooperativeKernel(test_kernel, dim3(tbs, 1, 1), dim3(1, 1, 1), args, max_smem,
                                    stream[i]);
    }
    cudaDeviceSynchronize();
    print_results<<<1, 1, 0, stream[0]>>>(bfr, ngrids * tbs);
    cudaDeviceSynchronize();
    return 0;
}

For every threadblock this prints the grid it belongs to, the blockIdx.x, the %smid, and the times as described above.

On H100, RTX5090, and RTX Pro 6000 (i.e. sm_90 and sm_120) this consistently shows all threadblocks in the first two grids launching at ~0ms, all blocks in the the third grid launching at ~1ms, and the fourth grid at ~2ms. This is exactly what I would expect (note the first grid runs for 9ms, the others for 1ms).

On RTX 5000 Ada, RTX A6000, and L40S I consistently see grids launched in a “split” fashion, with some threadblocks taking the remaining SMs while others have to wait for a previous grid to finish:

147 Grid  3, TB  0, SM 100 - Launched at 1004544ns, grid sync after 1001472ns, done waiting after 1001472ns, grid sync after 2000896ns
[42 lines omitted ...]
190 Grid  3, TB 43, SM  53 - Launched at 1004544ns, grid sync after 1001472ns, done waiting after 1001472ns, grid sync after 2000896ns
191 Grid  3, TB 44, SM  69 - Launched at 2004992ns, grid sync after 0ns, done waiting after 1000448ns, grid sync after 1000448ns
[3 lines omitted...]
195 Grid  3, TB 48, SM 123 - Launched at 2004992ns, grid sync after 0ns, done waiting after 1000448ns, grid sync after 1000448ns
4

Makes sense. Will this happen for all long-running kernels, including within the same process? Or will preemption only happen if multiple processes are competing for GPU resources? If I increase the wait time in my example to tens of seconds (needs some uint32_t→uint64_t fixes) and launch two simultaneous processes I can see preemption happening, but not when launching two grids on independent streams in the same process/CUDA context.

5

No it doesn’t apply to kernels launched from the same process. Context-switching (at least the way I am discussing it here) corresponds to inter-process activity, not intra-process.

6

OK I see the behavior you are describing (partial grid depositing) with your test case on L4 (cc8.9, CUDA 12.2). Although it is surprising to me, I don’t think it violates anything about cooperative grid behavior. That is, a cooperative grid is only guaranteed to work correctly when all blocks are scheduled (&). It is up to you to make sure that condition can be met at some point in your application trajectory. The fact that some blocks get scheduled “early” and others get scheduled when resources become available (later) doesn’t break anything, that I can see.

With respect to other processes that “you have no control over”, I have already pointed out that the interprocess case in the default condition provides for this via contex-switching (and for example if we were in the exclusive process compute mode “non-default” case, another process would be excluded.)

(As a single process/application example/thought experiment, let’s suppose you had a persistent kernel that runs “forever” (cooperative or not) and is using up most of the GPU, leaving 3 SM idle. Lets suppose at some later point you launch a cooperative kernel that requires more than these 3 SM – i.e. more than the resources available. This is a potential deadlock condition, whether we allow the 2nd cooperative launch to partially deposit on those 3 SMs, or not. We do hope that in the presence of many outstanding kernel launches, as resources become available, that the block scheduler/CWD will preferentially choose blocks from previously “started” cooperative launch(es).)

I will say that cooperative launches that are concurrent and requiring (in aggregate) not more than the full resources of the GPU doesn’t sound too scary to me. Likewise a cooperative launch with other non-cooperative throughput-style launches doesn’t sound scary. Cooperative launches that can’t be fully concurrent and also intending to use inter-kernel communication sounds very scary to me. It’s difficult to think through all the ramifications here, but inter-kernel communications is a frowned-upon design practice anyway.

(&) The documentation suggests:

To guarantee co-residency of the thread blocks on the GPU, the number of blocks launched needs to be carefully considered.

7

It’s good to know that preemption solves the problem across process boundaries. But as the kernel I’m writing is part of a library I don’t really have control over what other threads in the same process might be doing, either. I effectively have to look at each kernel launch in isolation. I was hoping that cooperative launches would offer some guarantee that all threadblocks will eventually be co-scheduled. If not, they seem to be no better than regular kernel launches.

Could you clarify what you mean by inter-kernel communication? I’m only communicating between threadblocks in the same grid, not across grids/launches, so this is no different from what cooperative_groups::this_grid().sync() does under the hood.

8

intra-kernel communication would be communicating between threads in the same kernel. CUDA recognizes this and has various mechanisms to facilitate. inter-kernel communication would be communicating between two concurrent kernels, running at the same time, where a thread in one kernel is somehow communicating with a thread in another kernel. CUDA provides no formal mechanisms that I can think of to help with this case (e.g. synchronization barriers, to pick just one example). I’m not including atomics as a “formal” mechanism, although you might argue that a device-wide semaphore (e.g. binary semaphore) might be a “formal” mechanism; I don’t know about that.

They provide a valid environment for the operation of grid sync, with all that that implies. An ordinary kernel launch does not. This is essentially the only salient feature of a cooperative launch, that I can think of.

Certainly CUDA does not provide any guarantees of preemption (for a single application, and its even a gray area in the multi-process case). Given that, I’m having a hard time coming up with an example of what an application calling your library could do that would have different (macroscopic) behavior in the wait-till-all-resources are ready case vs. the eager launch/deposit case. If the application has a long running kernel hogging resources, it will prevent completion of the cooperative kernel in either case.

Having said all that, for my own curiosity I might take a look at two things. 1. the granularity of globaltimer may be different on different GPU architectures. That might possibly be impacting observations here. 2. I might try a long-running resource hogging kernel to look at the deposit behavior, essentially similar to your test case. But these observations wouldn’t change anything I don’t think.

9

And a grid sync can only successfully complete if all threadblocks in the grid will eventually get co-scheduled. Which is exactly what I’m looking for. I think you’re right, the fact that some threadblocks start earlier than others isn’t an impediment to that. What I do think matters is that the remaining threadblocks of a cooperative launch get priority over other kernel launches enqueued on other streams. So I don’t end up with two or more partially scheduled grids that are deadlocked because neither can make forward progress if fewer than N threadblocks are running.

I’ve found the granularity of globaltimer to be between 32ns (sm_120) and 1024ns (sm_8x). The example uses millisecond scale intervals specifically so it won’t be affected by the granularity.

10

Yes, I think that is important, too. I mentioned it here:

I don’t have any knowledge of any such guarantee, but it seems to me logical that the CWD might behave that way. (it might already be an implicit guarantee in the existing mechanism.) If you think that is important, I would file a bug requesting doc clarification around that. I don’t know how it would be handled.