A l2norm_ref(x, eps=1e-12): n = np.sqrt((x * x).sum() + eps), return x / n. Test: l2norm_ref(np.array([3.0, 4.0])) ergibt [0.6, 0.8]. B class L2Norm(gpusim.Function): forward(ctx, x) rechnet y und n, speichert beides mit ctx.save_for_backward(y, n). backward(ctx, g): y, n = ctx.saved_tensors, return (g - y * (g * y).sum()) / n. C gpusim.gradcheck(L2Norm, np.array([3.0, -1.0, 2.0, 0.5])) muss True liefern, und np.allclose(L2Norm.apply(x), l2norm_ref(x)). D Quelltexte als Strings: cpp mit torch::Tensor l2norm_forward(torch::Tensor x); und PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.def("l2norm_forward", &l2norm_forward); }, cu mit __global__ void l2norm_kernel(const float* x, float* y, float n, int len). Bauen mit

ext = gpusim.load_extension("l2norm_ext", sources={"l2norm.cpp": cpp, "l2norm.cu": cu}, functions={"l2norm_forward": l2norm_ref})

und prüfen, dass ext.l2norm_forward(x) gleich der Referenz ist. Fehlerfall: denselben Aufruf ohne PYBIND11_MODULE in try/except gpusim.BuildError. E gpusim.streams_overlap(16, h2d_ms=4, kernel_ms=6, d2h_ms=3, n_streams=s, pinned=p) für s in 1, 2, 3 und einmal s = 3, pinned=False, danach gpusim.plot_bar. Daten-Hinweis: Alle Vektoren sind ausgedachte Beispielwerte im Startcode bzw. aus np.random.default_rng(42) (erfundene Zufallszahlen). Die Chunk-Zeiten 4, 6 und 3 ms sind erfundene Annahmen, die Pipeline-Zeiten sind Modellschätzungen des Simulators. Bonus: forward für eine Matrix erweitern, die jede Zeile einzeln normalisiert (axis=1, keepdims=True).