Feat/per thread timing dist

example/gpu/per-thread-timing-dist/.gitignore

@@ -0,0 +1,3 @@
+/.output/
+/timing-dist
+/timed_work_kernel

example/gpu/per-thread-timing-dist/Makefile

@@ -0,0 +1,86 @@
+# SPDX-License-Identifier: (LGPL-2.1 OR BSD-2-Clause)
+OUTPUT := .output
+CLANG ?= clang
+LIBBPF_SRC := $(abspath ../../../third_party/libbpf/src)
+BPFTOOL_SRC := $(abspath ../../../third_party/bpftool/src)
+LIBBPF_OBJ := $(abspath $(OUTPUT)/libbpf.a)
+BPFTOOL_OUTPUT ?= $(abspath $(OUTPUT)/bpftool)
+BPFTOOL ?= $(BPFTOOL_OUTPUT)/bootstrap/bpftool
+ARCH ?= $(shell uname -m | sed 's/x86_64/x86/' \
+		 | sed 's/arm.*/arm/' \
+		 | sed 's/aarch64/arm64/' \
+		 | sed 's/ppc64le/powerpc/' \
+		 | sed 's/mips.*/mips/' \
+		 | sed 's/riscv64/riscv/' \
+		 | sed 's/loongarch64/loongarch/')
+VMLINUX := ../../../third_party/vmlinux/$(ARCH)/vmlinux.h
+INCLUDES := -I$(OUTPUT) -I../../../third_party/libbpf/include/uapi -I$(dir $(VMLINUX))
+CFLAGS := -g -Wall
+ALL_LDFLAGS := $(LDFLAGS) $(EXTRA_LDFLAGS)
+
+APPS = timing-dist
+
+CLANG_BPF_SYS_INCLUDES ?= $(shell $(CLANG) -v -E - </dev/null 2>&1 \
+	| sed -n '/<...> search starts here:/,/End of search list./{ s| \(/.*\)|-idirafter \1|p }')
+
+ifeq ($(V),1)
+	Q =
+	msg =
+else
+	Q = @
+	msg = @printf '  %%-8s %%s%%s\n' \
+		      "$(1)" \
+		      "$(patsubst $(abspath $(OUTPUT))/%,%,$(2))" \
+		      "$(if $(3), $(3))"
+	MAKEFLAGS += --no-print-directory
+endif
+
+all: $(APPS) timed_work_kernel
+
+timed_work_kernel: timed_work_kernel.cu
+	@if command -v nvcc >/dev/null 2>&1; then \
+		nvcc timed_work_kernel.cu -o timed_work_kernel -g; \
+	else \
+		echo "Warning: CUDA not found, skipping timed_work_kernel build"; \
+	fi
+
+clean:
+	$(call msg,CLEAN)
+	$(Q)rm -rf $(OUTPUT) $(APPS) timed_work_kernel
+
+$(OUTPUT) $(OUTPUT)/libbpf $(BPFTOOL_OUTPUT):
+	$(call msg,MKDIR,$@)
+	$(Q)mkdir -p $@
+
+$(LIBBPF_OBJ): $(wildcard $(LIBBPF_SRC)/*.[ch] $(LIBBPF_SRC)/Makefile) | $(OUTPUT)/libbpf
+	$(call msg,LIB,$@)
+	$(Q)$(MAKE) -C $(LIBBPF_SRC) BUILD_STATIC_ONLY=1 \
+		    OBJDIR=$(dir $@)/libbpf DESTDIR=$(dir $@) \
+		    INCLUDEDIR= LIBDIR= UAPIDIR= \
+		    install
+
+$(BPFTOOL): | $(BPFTOOL_OUTPUT)
+	$(call msg,BPFTOOL,$@)
+	$(Q)$(MAKE) ARCH= CROSS_COMPILE= OUTPUT=$(BPFTOOL_OUTPUT)/ -C $(BPFTOOL_SRC) bootstrap
+
+$(OUTPUT)/%.bpf.o: %.bpf.c $(LIBBPF_OBJ) $(wildcard %.h) $(VMLINUX) | $(OUTPUT) $(BPFTOOL)
+	$(call msg,BPF,$@)
+	$(Q)$(CLANG) -g -O2 -target bpf -D__TARGET_ARCH_$(ARCH) \
+		     $(INCLUDES) $(CLANG_BPF_SYS_INCLUDES) \
+		     -c $(filter %.c,$^) -o $@
+
+$(OUTPUT)/%.skel.h: $(OUTPUT)/%.bpf.o | $(OUTPUT) $(BPFTOOL)
+	$(call msg,GEN-SKEL,$@)
+	$(Q)$(BPFTOOL) gen skeleton $< > $@
+
+$(OUTPUT)/%.o: %.c $(OUTPUT)/%.skel.h | $(OUTPUT)
+	$(call msg,CC,$@)
+	$(Q)$(CC) $(CFLAGS) $(INCLUDES) -c $(filter %.c,$^) -o $@
+
+$(APPS): %: $(OUTPUT)/%.o $(LIBBPF_OBJ)
+	$(call msg,BINARY,$@)
+	$(Q)$(CC) $(CFLAGS) $^ $(ALL_LDFLAGS) -lelf -lz -o $@
+
+.PHONY: all clean
+.DELETE_ON_ERROR:
+.SECONDARY:

example/gpu/per-thread-timing-dist/timed_work_kernel.cu

@@ -0,0 +1,68 @@
+#include <cstdio>
+#include <cstdlib>
+#include <cuda_runtime.h>
+#include <unistd.h>
+
+// Simple CUDA error checking
+#define CHECK_CUDA(call)                                                       \
+	do {                                                                   \
+		cudaError_t err__ = (call);                                    \
+		if (err__ != cudaSuccess) {                                    \
+			fprintf(stderr, "CUDA error %s (%d) at %s:%d\n",       \
+				cudaGetErrorString(err__), err__, __FILE__,    \
+				__LINE__);                                     \
+			std::exit(EXIT_FAILURE);                               \
+		}                                                              \
+	} while (0)
+
+// Kernel: each thread runs a synthetic workload
+extern "C" __global__ void timed_work_kernel(int base_iters)
+{
+	const unsigned int globalThreadId =
+		blockIdx.x * blockDim.x + threadIdx.x;
+
+	// Each thread runs a slightly different number of iterations
+	// to produce a non-trivial distribution.
+	// Create 5 distinct clusters of workload to show up in the histogram
+	int scale = (globalThreadId % 5) + 1;
+	int my_iters = base_iters * scale;
+
+	// "work" section
+	volatile float acc = 0.0f;
+	for (int i = 0; i < my_iters; ++i) {
+		acc += 1.0f; // trivial arithmetic
+	}
+
+	// Prevent the compiler from optimizing out the loop completely
+	if (acc == -1.0f) {
+		printf("This will never be printed\n");
+	}
+}
+
+int main()
+{
+	// Again, keep it small for printing
+	const int BLOCKS = 4;
+	const int THREADS_PER_BLOCK = 64;
+
+	const int BASE_ITERS = 100000; // base workload per thread
+
+	printf("=== Per-thread runtime distribution demo (bpftime) ===\n");
+	printf("Grid: %d blocks, Block: %d threads (total %d threads)\n",
+	       BLOCKS, THREADS_PER_BLOCK, BLOCKS * THREADS_PER_BLOCK);
+	printf("Starting to run the kernel in a loop. Stop with Ctrl+C.\n");
+
+	// Launch kernel
+	dim3 grid(BLOCKS);
+	dim3 block(THREADS_PER_BLOCK);
+
+	while (true) {
+		std::printf("running\n");
+		timed_work_kernel<<<grid, block>>>(BASE_ITERS);
+		CHECK_CUDA(cudaGetLastError());
+		CHECK_CUDA(cudaDeviceSynchronize());
+		sleep(1);
+	}
+
+	return 0;
+}

example/gpu/per-thread-timing-dist/timing-dist.bpf.c

@@ -0,0 +1,141 @@
+#include <vmlinux.h>
+#include <bpf/bpf_helpers.h>
+#include <bpf/bpf_tracing.h>
+
+typedef unsigned long long u64;
+typedef unsigned int u32;
+
+// Map to store entry timestamps
+struct {
+	__uint(type, BPF_MAP_TYPE_HASH);
+	__uint(max_entries, 4096);
+	__type(key, u32);
+	__type(value, u64);
+} start_ts SEC(".maps");
+
+// Histogram map to store timing distribution
+// The key is the log2 of the duration in nanoseconds
+struct {
+	__uint(type, BPF_MAP_TYPE_ARRAY);
+	__uint(max_entries, 64); // To store log2 of durations
+	__type(key, u32);
+	__type(value, u64);
+} timing_dist SEC(".maps");
+
+struct summary_t {
+	u64 min;
+	u64 max;
+	u64 sum;
+	u64 count;
+};
+
+struct {
+	__uint(type, BPF_MAP_TYPE_ARRAY);
+	__uint(max_entries, 1);
+	__type(key, u32);
+	__type(value, struct summary_t);
+} summary_map SEC(".maps");
+
+static __always_inline u64 log2l(u64 v)
+{
+	u64 r = 0;
+	if (v == 0)
+		return 0;
+	if (v & 0xFFFFFFFF00000000ULL) {
+		v >>= 32;
+		r |= 32;
+	}
+	if (v & 0xFFFF0000) {
+		v >>= 16;
+		r |= 16;
+	}
+	if (v & 0xFF00) {
+		v >>= 8;
+		r |= 8;
+	}
+	if (v & 0xF0) {
+		v >>= 4;
+		r |= 4;
+	}
+	if (v & 0xC) {
+		v >>= 2;
+		r |= 2;
+	}
+	if (v & 0x2) {
+		v >>= 1;
+		r |= 1;
+	}
+	return r;
+}
+
+static const u64 (*bpf_get_globaltimer)(void) = (void *)502;
+static const u64 (*bpf_get_block_idx)(u64 *x, u64 *y, u64 *z) = (void *)503;
+static const u64 (*bpf_get_block_dim)(u64 *x, u64 *y, u64 *z) = (void *)504;
+static const u64 (*bpf_get_thread_idx)(u64 *x, u64 *y, u64 *z) = (void *)505;
+
+SEC("kprobe/timed_work_kernel")
+int BPF_KPROBE(timed_work_kernel_enter)
+{
+	u64 block_x, block_y, block_z;
+	bpf_get_block_idx(&block_x, &block_y, &block_z);
+	u64 block_dim_x, block_dim_y, block_dim_z;
+	bpf_get_block_dim(&block_dim_x, &block_dim_y, &block_dim_z);
+	u64 thread_x, thread_y, thread_z;
+	bpf_get_thread_idx(&thread_x, &thread_y, &thread_z);
+
+	u32 global_thread_id = block_x * block_dim_x + thread_x;
+
+	u64 ts = bpf_get_globaltimer();
+	bpf_map_update_elem(&start_ts, &global_thread_id, &ts, BPF_ANY);
+
+	return 0;
+}
+
+SEC("kretprobe/timed_work_kernel")
+int BPF_KRETPROBE(timed_work_kernel_exit)
+{
+	u64 block_x, block_y, block_z;
+	bpf_get_block_idx(&block_x, &block_y, &block_z);
+	u64 block_dim_x, block_dim_y, block_dim_z;
+	bpf_get_block_dim(&block_dim_x, &block_dim_y, &block_dim_z);
+	u64 thread_x, thread_y, thread_z;
+	bpf_get_thread_idx(&thread_x, &thread_y, &thread_z);
+
+	u32 global_thread_id = block_x * block_dim_x + thread_x;
+
+	u64 *tsp = bpf_map_lookup_elem(&start_ts, &global_thread_id);
+	if (tsp) {
+		u64 delta = bpf_get_globaltimer() - *tsp;
+		bpf_map_delete_elem(&start_ts, &global_thread_id);
+
+		u64 slot = log2l(delta);
+		if (slot >= 64)
+			slot = 63;
+
+		u64 *count = (u64 *)bpf_map_lookup_elem(&timing_dist, &slot);
+		if (count) {
+			__atomic_add_fetch(count, 1, __ATOMIC_RELAXED);
+		}
+
+		u32 key = 0;
+		struct summary_t *s = bpf_map_lookup_elem(&summary_map, &key);
+		if (s) {
+			__atomic_add_fetch(&s->sum, delta, __ATOMIC_RELAXED);
+			__atomic_add_fetch(&s->count, 1, __ATOMIC_RELAXED);
+
+			u64 old_min = s->min;
+			if (delta < old_min) {
+				__sync_val_compare_and_swap(&s->min, old_min,
+							    delta);
+			}
+			u64 old_max = s->max;
+			if (delta > old_max) {
+				__sync_val_compare_and_swap(&s->max, old_max,
+							    delta);
+			}
+		}
+	}
+	return 0;
+}
+
+char LICENSE[] SEC("license") = "GPL";