feat: support early exit (#459)

[Forsworns]

Description

• Add two bpf helper functions for CUDA.
• Add an early-exit demo in CUDA examples.
  Fixes [FEATURE] CUDA kernel early exit demo #459

Type of change

• New feature (non-breaking change which adds functionality)

How Has This Been Tested?

Tested via the provided demo.

[Forsworns]

I just found there were some typos in the README and the binary was included.

[yunwei37]

Thanks a lot!

[Copilot]

Pull Request Overview

This PR adds early exit functionality for CUDA kernels through eBPF helpers, enabling kernel atomization capabilities similar to network packet filtering. It introduces two new BPF helper functions (bpf_cuda_exit and bpf_get_grid_dim) and provides a complete demonstration through a vector addition example with partition-based execution control.

Key changes:

• Two new BPF helper functions (507: exit, 508: get_grid_dim) for CUDA kernel control
• Complete atomizer example with partition-based block filtering
• PTX-level early exit implementation via inline assembly

Reviewed Changes

Copilot reviewed 11 out of 12 changed files in this pull request and generated 6 comments.

Show a summary per file

File	Description
attach/nv_attach_impl/trampoline/default_trampoline.cu	Implements the two new BPF helper functions in CUDA
attach/nv_attach_impl/trampoline_ptx.h	Adds PTX assembly for the new helper functions
attach/nv_attach_impl/nv_attach_impl_patcher.cpp	Registers the new helper functions (507 and 508)
example/gpu/atomizer/atomizer.bpf.c	eBPF program implementing partition-based kernel atomization
example/gpu/atomizer/atomizer.c	Userspace loader for the eBPF atomizer program
example/gpu/atomizer/vec_add.cu	CUDA vector addition demo application
example/gpu/atomizer/main.ptx	Generated PTX assembly from the demo
example/gpu/atomizer/filter_hashtag.py	Utility script to filter preprocessor directives
example/gpu/atomizer/README.md	Documentation for the atomizer example
example/gpu/atomizer/Makefile	Build configuration for the atomizer example
example/gpu/atomizer/.gitignore	Git ignore rules for build artifacts

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[Forsworns]

I have addressed typos. But I found another problem:

when I set the launching configuration in vector_add.cu to vectorAdd<<<10, 1>>>(d_A, d_B, d_C);,
and try to read the pre-configured partition number/index from the BPF maps in the atomizer.bpf.c, it sometimes returns null pointer for the given key. I'm not sure where the problem is.

Currently, I only launch a single block in the vector_add.c and it works well. But then only part of threads are exited, instead of the whole thread blocks. Thus, the semantic is different from the LithOS. :(

[yunwei37]

I think we need to provide new attach types instead of function probes to support that semantic. That's not hard as we have seperated the attach types into passes?

[yunwei37]

Maybe we can merge that and continue future work on next PR? @Officeyutong @Sy0307 @Forsworns

[Forsworns]

Maybe we can merge that and continue future work on next PR? @Officeyutong @Sy0307 @Forsworns

I'm fine about the attach type. But I'm still bothered by the above BPF map issue, do you have any ideas?

[Sy0307]

Maybe I can review it later for BPF map issue? Can you give a more detailed description?

[Forsworns]

Maybe I can review it later for BPF map issue? Can you give a more detailed description?

@Sy0307 I just opened a new issue in #486 and provide two examples. I guess it is related to the synchronize between host and device.