fix(ncu_profile): auto-add __global__ to kernels that omit it

Some .cu files declare "extern \"C\" void solve(...)" without __global__.
Without __global__, nvcc compiles the function as a host function and the
triple-chevron launch syntax fails. Add _ensure_global() preprocessing
(mirrors benchmark.py) to insert __global__ when missing.

Co-Authored-By: Claude <noreply@anthropic.com>

Author: grissom.wang

examples/vectorAdd/solution_bench_wrapper.cu

@@ -0,0 +1,67 @@
+// Auto-generated NCU profiling bench — self-contained, no subprocess.
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <cuda_runtime.h>
+
+// User kernel source
+#include "/Users/wangtianqing/Project/skills/cuda-code-skill/examples/vectorAdd/solution.cu"
+
+int main(int argc, char **argv) {
+    int warmup = 10, repeat = 22;
+        int N = 1000000;
+        
+        for (int i = 1; i < argc; i++) {
+            if (strncmp(argv[i], "--N=", 4) == 0) N = atoi(argv[i] + 4);
+            if (strncmp(argv[i], "--warmup=", 9) == 0) warmup = atoi(argv[i] + 9);
+            if (strncmp(argv[i], "--repeat=", 9) == 0) repeat = atoi(argv[i] + 9);
+        }
+
+    int threads = 256, blocks = (N + 255) / 256;
+
+    // Allocate device buffers
+    size_t bytes_A = (size_t)N * 4;
+        float *d_A;
+        cudaMalloc(&d_A, bytes_A);
+        size_t bytes_B = (size_t)N * 4;
+        float *d_B;
+        cudaMalloc(&d_B, bytes_B);
+        size_t bytes_C = (size_t)N * 4;
+        float *d_C;
+        cudaMalloc(&d_C, bytes_C);
+
+    // Initialize: non-zero pattern for inputs, zero for outputs
+    cudaMemset(d_A, 205, bytes_A);
+        cudaMemset(d_B, 205, bytes_B);
+        cudaMemset(d_C, 0, bytes_C);
+
+    cudaDeviceSynchronize();
+
+    // Warmup
+    for (int i = 0; i < warmup; i++)
+        solve<<<blocks, threads>>>(d_A, d_B, d_C, N);
+    cudaDeviceSynchronize();
+
+    // Timed iterations
+    cudaEvent_t start, stop;
+    cudaEventCreate(&start); cudaEventCreate(&stop);
+    double total = 0.0;
+    for (int i = 0; i < repeat; i++) {
+        cudaEventRecord(start, 0);
+        solve<<<blocks, threads>>>(d_A, d_B, d_C, N);
+        cudaEventRecord(stop, 0);
+        cudaEventSynchronize(stop);
+        float ms; cudaEventElapsedTime(&ms, start, stop);
+        total += ms;
+        if (i < 2 || i >= repeat - 2) printf("  iter %d: %.4f ms\n", i, ms);
+    }
+    printf("  avg: %.4f ms\n", total / repeat);
+
+    // Cleanup
+    cudaEventDestroy(start); cudaEventDestroy(stop);
+    cudaFree(d_A);
+    cudaFree(d_B);
+    cudaFree(d_C);
+    printf("done\n");
+    return 0;
+}

skills/kernel-benchmarker/scripts/ncu_profile.py

@@ -96,6 +96,28 @@ def parse_solve_signature(cu_file: str):
     return params
 
 
+def _ensure_global(cu_file: str) -> str:
+    """If solve lacks __global__, return path to a temp copy with __global__ added."""
+    with open(cu_file, "r") as f:
+        src = f.read()
+
+    if re.search(r'extern\s+"C"\s+__global__\s+void\s+solve', src):
+        return cu_file
+
+    new_src = re.sub(
+        r'(extern\s+"C"\s+)void\s+solve',
+        r'\1__global__ void solve',
+        src
+    )
+    if new_src == src:
+        return cu_file
+
+    tmp = cu_file + ".global.cu"
+    with open(tmp, "w") as f:
+        f.write(new_src)
+    return tmp
+
+
 def detect_arch() -> str:
     try:
         import torch
@@ -289,10 +311,13 @@ def main():
             except ValueError:
                 pass
 
-    cu_abs = os.path.abspath(args.cu_file)
+    # Preprocess: add __global__ if missing (some kernels omit it)
+    cu_file = _ensure_global(args.cu_file)
+    cu_abs = os.path.abspath(cu_file)
     arch = args.arch or detect_arch()
 
-    # Parse kernel signature
+    # Parse kernel signature (from original file since _ensure_global
+    # only adds __global__, doesn't change the parameter list)
     try:
         params = parse_solve_signature(args.cu_file)
     except ValueError as e:
@@ -323,6 +348,9 @@ def main():
     print(f"[compile] {' '.join(cmd)}", flush=True)
     result = subprocess.run(cmd, capture_output=True, text=True)
     os.remove(wrapper_path)
+    # Clean up temp file from _ensure_global if one was created
+    if cu_file != args.cu_file and os.path.exists(cu_file):
+        os.remove(cu_file)
 
     if result.returncode != 0:
         print(f"Compilation failed:\n{result.stderr}", file=sys.stderr)