Skill v1.0.1

currentAutomated scan100/100

gpu-mode/popcorn-cli/load-inline-native-code

3 files

──Details

PublishedOctober 25, 2026 at 03:01 PM

Content Hashsha256:fb9e60c9c36b1d9c...

Git SHA036b3ead604f

Bump Typepatch

Compare with v1.0.0

──Files

Files (1 file, 4.6 KB)

SKILL.md4.6 KBactive

SKILL.md · 145 lines · 4.6 KB

version: "1.0.1" name: load-inline-native-code description: Helps write CUDA and HIP kernels using torch.utils.cpp_extension.load_inline(). Use when users want to write native GPU code (CUDA/HIP) inside a Python submission file. compatibility: Intended for popcorn-cli submissions targeting NVIDIA or AMD GPUs with native kernel code.

Writing Native GPU Kernels with load_inline()

Use this skill when the user wants to write a custom CUDA or HIP kernel inside their Python submission file using torch.utils.cpp_extension.load_inline().

Overview

load_inline() compiles C++/CUDA/HIP source code at runtime and loads it as a Python module. This lets you write raw GPU kernels directly in your submission.py without a separate build system.

CUDA Template (NVIDIA GPUs)

python

import torch
from torch.utils.cpp_extension import load_inline
from task import input_t, output_t
CUDA_SRC = """
template <typename scalar_t>
__global__ void my_kernel(const scalar_t* __restrict__ input,
                          scalar_t* __restrict__ output,
                          int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        output[idx] = input[idx];
    }
}
torch::Tensor my_op(torch::Tensor input, torch::Tensor output) {
    int N = input.numel();
    const int threads = 256;
    const int blocks = (N + threads - 1) / threads;
    AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.scalar_type(), "my_kernel", ([&] {
        my_kernel<scalar_t><<<blocks, threads>>>(
            input.data_ptr<scalar_t>(),
            output.data_ptr<scalar_t>(),
            N
        );
    }));
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        throw std::runtime_error(cudaGetErrorString(err));
    }
    return output;
}
"""
CPP_SRC = """
torch::Tensor my_op(torch::Tensor input, torch::Tensor output);
"""
module = load_inline(
    name='my_module',
    cpp_sources=[CPP_SRC],
    cuda_sources=[CUDA_SRC],
    functions=['my_op'],
    verbose=True,
)
def custom_kernel(data: input_t) -> output_t:
    input, output = data
    return module.my_op(input, output)

HIP Template (AMD GPUs)

python

import os
os.environ['PYTORCH_ROCM_ARCH'] = 'gfx942'
os.environ['CXX'] = 'clang++'
import torch
from torch.utils.cpp_extension import load_inline
from task import input_t, output_t
CUDA_SRC = """
#include <hip/amd_detail/amd_hip_bf16.h>
__global__ void my_kernel(const float* input, float* output, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        output[idx] = input[idx];
    }
}
void my_op(torch::Tensor input, torch::Tensor output) {
    int N = input.numel();
    const int threads = 256;
    const int blocks = (N + threads - 1) / threads;
    my_kernel<<<blocks, threads>>>(
        input.data_ptr<float>(),
        output.data_ptr<float>(),
        N
    );
}
"""
CPP_SRC = """
void my_op(torch::Tensor input, torch::Tensor output);
"""
module = load_inline(
    name='my_module',
    cpp_sources=[CPP_SRC],
    cuda_sources=[CUDA_SRC],
    functions=['my_op'],
    verbose=True,
    extra_cuda_cflags=["--offload-arch=gfx942", "-std=c++20"],
)
def custom_kernel(data: input_t) -> output_t:
    input, output = data
    module.my_op(input, output)
    return output

Key Points

cpp_sources: C++ header declaring the functions you want to call from Python. These are the bindings.
cuda_sources: The actual CUDA/HIP kernel code and the C++ wrapper that launches it.
functions: List of function names to expose to Python. Must match the C++ function signatures exactly.
verbose=True: Prints compilation output so you can debug build errors.
extra_cuda_cflags: Pass extra compiler flags. Needed for AMD HIP (--offload-arch=gfx942) or C++ standard selection.

Common Patterns

Use AT_DISPATCH_FLOATING_TYPES_AND_HALF to handle multiple dtypes in CUDA kernels.
For AMD/HIP, set PYTORCH_ROCM_ARCH and CXX env vars before importing torch.
Always check cudaGetLastError() after kernel launches for NVIDIA targets.
The load_inline call compiles on first run and caches the result. Subsequent runs reuse the cache unless the source changes.
Keep the module-level load_inline() call outside custom_kernel() so compilation happens once at import time, not on every call.

Guardrails

The custom_kernel function signature must match def custom_kernel(data: input_t) -> output_t:.
The module is compiled at import time. Do not call load_inline() inside custom_kernel().
For AMD GPUs, always set PYTORCH_ROCM_ARCH before any torch import.
Use torch::Tensor in C++ signatures for seamless Python-C++ tensor passing.

← v1.0.0 All versions