gpt-5.5-priority · popcorn · l12_default_gpt
correct 22_SE3InvariantLinear.py
model
gpt-5.5-priority
turns / max
9 / 10
tool calls
9
tokens (total)
0
wall clock
—
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.02 μs
ref runtime
0.10 μs
speedup
4.48x
started
2026-05-02T17:49:15Z
timingkernelbench.timing
kernel runtime
0.02 μs
ref runtime
0.10 μs
speedup
4.48x
kernel mean / median / std
0.02 / — / 0.01 μs
ref mean / median / std
0.10 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
max abs error
7.15e-07
mean abs error
4.15e-08
max rel error
0.05
mean rel error
5.44e-07
gpu memoryextended_metrics.measure_memory
kernel peak
33.93 MB
ref peak
34.55 MB
kernel / ref
0.98x
kernel launches / fusiontorch.profiler
kernel launches
2
ref launches
51
fusion ratio (ref / kernel)
25.50x
total cuda time
33 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.008
dram util %
0.8%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
26.59 GB/s
dram util %
0.8%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
39.8%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
88.6%
ld sectors / req
24.00
st sectors / req
—
regs / thread
37
smem / block
1664 B
block size
128
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
mio_throttle
3476.1%
short_scoreboard
2078.2%
long_scoreboard
831.9%
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    SE(3)-invariant linear layer.  Maps node features that are augmented
    with 3-D coordinate information through a linear transform that
    respects rotational and translational invariance.  The layer operates
    on (scalar_features, vector_features) pairs where scalar features
    transform trivially and vector features transform as 3-D vectors.
    Used in equivariant GNNs (EGNN, PaiNN, TFN).
    """

    def __init__(self, scalar_in, scalar_out, vector_in, vector_out):
        super().__init__()
        self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)
        self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))
        nn.init.xavier_uniform_(self.vector_weight)
        self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)

    def forward(
        self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor
    ) -> tuple:
        """
        Args:
            scalar_feat: (B, N, scalar_in) – invariant node features
            vector_feat: (B, N, vector_in, 3) – equivariant vector features
        Returns:
            (scalar_out, vector_out):
                scalar_out: (B, N, scalar_out)
                vector_out: (B, N, vector_out, 3)
        """
        # Invariant scalars from vector norms
        vec_norm = vector_feat.norm(dim=-1)  # (B, N, vector_in)
        combined = torch.cat([scalar_feat, vec_norm], dim=-1)

        s_out = self.scalar_linear(combined)

        # Equivariant vector transform: linear combination of input vectors
        v_out = torch.einsum("oi,bnid->bnod", self.vector_weight, vector_feat)
        # Gating by scalar-derived signal
        gate = torch.sigmoid(self.vector_gate(combined)).unsqueeze(-1)  # (B, N, vector_out, 1)
        v_out = gate * v_out

        return s_out, v_out


scalar_in = 64
scalar_out = 64
vector_in = 16
vector_out = 16
num_nodes = 128
batch_size = 8


def get_inputs():
    p = popcorn_pri
    bs = p.jitter_int(batch_size)
    nn = p.jitter_int(num_nodes)
    s = torch.randn(bs, nn, scalar_in)
    v = torch.randn(bs, nn, vector_in, 3)
    return [s, v]


def get_init_inputs():
    return [scalar_in, scalar_out, vector_in, vector_out]
turn 013.3s turn · 13.3s LLM · 0.0s tools
reasoning
**Planning CUDA extension** I need to craft a Python file with a CUDA extension. It might be important to know the GPU specs for this. I think I should define the model parameters and ensure they match inherited settings, using ModelNew(nn.Module) with the same initialization for nn.Linear. I need to consider custom forward functions and fixed dimensions. Optimizing for fixed sizes might be useful, and I'll need to compile inline to manage computations efficiently. I should explore the references for scalar and vector operations in CUDA. I'll consider kernels and likely stay within small memory usage to avoid PyTorch overhead. **Designing CUDA thread structure** I’m planning for one block per BN, maybe using 128 or 256 threads. I think I can use shared memory for scalar outputs, norms, and perhaps vector features. Each thread might compute one scalar output and handle gate and vector sums. With 1024 blocks, the overhead seems acceptable. Each block loads combined outputs repeatedly. I see the potential for using one block with 128 threads, where threads 0-63 handle scalar outputs, threads 64-79 manage gate pre-computations, and threads 80-127 could work on vector outputs. I should also explore shared memory caching to improve load efficiency. **Addressing occupancy and latency** I'm considering the small occupancy of the blocks and the potential latency issues that might arise. I think I need to tackle the binding, which is essential for efficient operations. To facilitate this, I plan to use a C++ function that returns a `vector<torch::Tensor>`. This approach seems like a solid way to manage the tensors effectively while optimizing for performance. I just have to ensure that everything integrates smoothly!
call → get_gpu_specs
callget_gpu_specs {}
result ← get_gpu_specs (ok)
okget_gpu_specs GPU specs for NVIDIA H100 80GB HBM3: total memory (runtime): 79.1 GB spec entry: H100 GPU Architecture: Hopper GPU Memory: 80GB Memory Bandwidth: 3.35 TB/s FP64 TFLOPS: 34 FP64 Tensor Core TFLOPS: 67 FP32 TFLOPS: 67 TF32 Tensor Core TFLOPS: 989 with sparsity BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity FP16 Tensor Core TFLOPS: 1979 with sparsity FP8 Tensor Core TFLOPS: 3958 with sparsity INT8 Tensor Core TOPS: 3958 with sparsity Register File Size: 64K 32-bit registers per SM Maximum number of registers per thread: 255 Maximum number of thread blocks per SM: 32 Shared memory capacity per SM: 228 KB Maximum shared memory per thread block: 227 KB
turn 1 · 1 tool call(s) · 29.2s turn · 29.2s LLM · 8,946 chars (click to expand)
turn 129.2s turn · 29.2s LLM · 0.0s tools
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out,\n int rows, int scalar_in, int scalar_out_dim, int vector_in, int vector_out_dim)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n int C = scalar_in + vector_in;\n extern __shared__ float sh[];\n float* comb = sh; // C\n float* gate = sh + C; // vector_out_dim\n float* vec = gate + vector_out_dim; // vector_in*3\n\n const float* srow = scalar_feat + (size_t)row * scalar_in;\n const float* vrow = vector_feat + (size_t)row * vector_in * 3;\n\n for (int i = tid; i < scalar_in; i += blockDim.x) {\n comb[i] = srow[i];\n }\n for (int i = tid; i < vector_in * 3; i += blockDim.x) {\n vec[i] = vrow[i];\n }\n __syncthreads();\n\n for (int i = tid; i < vector_in; i += blockDim.x) {\n float x = vec[i*3 + 0];\n float y = vec[i*3 + 1];\n float z = vec[i*3 + 2];\n comb[scalar_in + i] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n for (int o = tid; o < scalar_out_dim; o += blockDim.x) {\n const float* w = scalar_w + (size_t)o * C;\n float acc = scalar_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) {\n acc = fmaf(w[i], comb[i], acc);\n }\n scalar_out[(size_t)row * scalar_out_dim + o] = acc;\n }\n\n for (int o = tid; o < vector_out_dim; o += blockDim.x) {\n const float* w = gate_w + (size_t)o * C;\n float acc = gate_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) {\n acc = fmaf(w[i], comb[i], acc);\n }\n gate[o] = 1.0f / (1.0f + expf(-acc));\n }\n __syncthreads();\n\n int total_v = vector_out_dim * 3;\n for (int idx = tid; idx < total_v; idx += blockDim.x) {\n int o = idx / 3;\n int d = idx - o * 3;\n const float* w = vector_w + (size_t)o * vector_in;\n float acc = 0.0f;\n #pragma unroll 4\n for (int i = 0; i < vector_in; ++i) {\n acc = fmaf(w[i], vec[i*3 + d], acc);\n }\n vector_out[((size_t)row * vector_out_dim + o) * 3 + d] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n int scalar_in = (int)scalar_feat.size(2);\n int vector_in = (int)vector_feat.size(2);\n int scalar_out_dim = (int)scalar_w.size(0);\n int vector_out_dim = (int)vector_w.size(0);\n auto s_out = torch::empty({B, N, scalar_out_dim}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_out_dim, 3}, vector_feat.options());\n int rows = (int)(B * N);\n int threads = 128;\n size_t shmem = (size_t)(scalar_in + vector_in + vector_out_dim + vector_in * 3) * sizeof(float);\n se3_forward_kernel<<<rows, threads, shmem, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(),\n vector_w.data_ptr<float>(), gate_w.data_ptr<float>(), gate_b.data_ptr<float>(),\n s_out.data_ptr<float>(), v_out.data_ptr<float>(),\n rows, scalar_in, scalar_out_dim, vector_in, vector_out_dim);\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(\n name='se3_fused_ext_v1',\n cpp_sources='',\n cuda_sources=_cuda_src,\n functions=['se3_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat,\n self.scalar_linear.weight, self.scalar_linear.bias,\n self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: RuntimeError. RuntimeError: Error building extension 'se3_fused_ext_v1': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=se3_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/tejas/PopcornBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/tejas/PopcornBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/include/python3.10 -fPIC -std=c++17 -O3 -c /scratch/tejas/PopcornBench/runs/l12_default/popcorn/gpt-5.5-priority/level_1_problem_22_cache/k_866448e6bfe2/se3_fused_ext_v1/main.cpp -o main.o FAILED: [code=1] main.o /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=se3_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/tejas/PopcornBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/tejas/PopcornBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/include/python3.10 -fPIC -std=c++17 -O3 -c /scratch/tejas/PopcornBench/runs/l12_default/popcorn/gpt-5.5-priority/level_1_problem_22_cache/k_866448e6bfe2/se3_fused_ext_v1/main.cpp -o main.o /scratch/tejas/PopcornBench/runs/l12_default/popcorn/gpt-5.5-priority/level_1_problem_22_cache/k_866448e6bfe2/se3_fused_ext_v1/main.cpp: In function 'void pybind11_init_se3_fused_ext_v1(pybind11::module_&)': /scratch/tejas/PopcornBench/runs/l12_default/popcorn/gpt-5.5-priority/level_1_problem_22_cache/k_866448e6bfe2/se3_fused_ext_v1/main.cpp:4:50: error: 'se3_forward' was not declared in this scope 4 | m.def("se3_forward", torch::wrap_pybind_function(se3_forward), "se3_forward"); | ^~~~~~~~~~~ [2/3] /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=se3_fused_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/tejas/PopcornBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/tejas/PopcornBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /home/prabhune/.local/share/uv/python/cpython-3.10-linux-x86_64-gnu/include/python3.10 -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_90,code=sm_90 --compiler-options '-fPIC' -O3 --use_fast_math -std=c++17 -c /scratch/tejas/PopcornBench/runs/l12_default/popcorn/gpt-5.5-priority/level_1_problem_22_cache/k_866448e6bfe2/se3_fused_ext_v1/cuda.cu -o cuda.cuda.o nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used ninja: build stopped: subcommand failed.
turn 2 · 1 tool call(s) · 12.0s turn · 12.0s LLM · 6,498 chars (click to expand)
turn 212.0s turn · 12.0s LLM · 0.0s tools
reasoning
**Organizing function declarations** I see that I need to declare functions in the cpp_sources for the auto main cpp. To make sure everything works, I should also include vector and torch. It looks like I’ll use the cuda_sources definition for that part. Let's make sure I get the right structure so that the code runs smoothly. Keeping things organized is key to avoiding issues later!
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out,\n int rows, int scalar_in, int scalar_out_dim, int vector_in, int vector_out_dim)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n int C = scalar_in + vector_in;\n extern __shared__ float sh[];\n float* comb = sh;\n float* gate = sh + C;\n float* vec = gate + vector_out_dim;\n\n const float* srow = scalar_feat + (size_t)row * scalar_in;\n const float* vrow = vector_feat + (size_t)row * vector_in * 3;\n\n for (int i = tid; i < scalar_in; i += blockDim.x) comb[i] = srow[i];\n for (int i = tid; i < vector_in * 3; i += blockDim.x) vec[i] = vrow[i];\n __syncthreads();\n\n for (int i = tid; i < vector_in; i += blockDim.x) {\n float x = vec[i*3 + 0], y = vec[i*3 + 1], z = vec[i*3 + 2];\n comb[scalar_in + i] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n for (int o = tid; o < scalar_out_dim; o += blockDim.x) {\n const float* w = scalar_w + (size_t)o * C;\n float acc = scalar_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) acc = fmaf(w[i], comb[i], acc);\n scalar_out[(size_t)row * scalar_out_dim + o] = acc;\n }\n\n for (int o = tid; o < vector_out_dim; o += blockDim.x) {\n const float* w = gate_w + (size_t)o * C;\n float acc = gate_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) acc = fmaf(w[i], comb[i], acc);\n gate[o] = 1.0f / (1.0f + expf(-acc));\n }\n __syncthreads();\n\n int total_v = vector_out_dim * 3;\n for (int idx = tid; idx < total_v; idx += blockDim.x) {\n int o = idx / 3;\n int d = idx - o * 3;\n const float* w = vector_w + (size_t)o * vector_in;\n float acc = 0.0f;\n #pragma unroll 4\n for (int i = 0; i < vector_in; ++i) acc = fmaf(w[i], vec[i*3 + d], acc);\n vector_out[((size_t)row * vector_out_dim + o) * 3 + d] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n int scalar_in = (int)scalar_feat.size(2);\n int vector_in = (int)vector_feat.size(2);\n int scalar_out_dim = (int)scalar_w.size(0);\n int vector_out_dim = (int)vector_w.size(0);\n auto s_out = torch::empty({B, N, scalar_out_dim}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_out_dim, 3}, vector_feat.options());\n int rows = (int)(B * N);\n int threads = 128;\n size_t shmem = (size_t)(scalar_in + vector_in + vector_out_dim + vector_in * 3) * sizeof(float);\n se3_forward_kernel<<<rows, threads, shmem, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(),\n vector_w.data_ptr<float>(), gate_w.data_ptr<float>(), gate_b.data_ptr<float>(),\n s_out.data_ptr<float>(), v_out.data_ptr<float>(),\n rows, scalar_in, scalar_out_dim, vector_in, vector_out_dim);\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(\n name='se3_fused_ext_v2',\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=['se3_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat,\n self.scalar_linear.weight, self.scalar_linear.bias,\n self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 3 · 1 tool call(s) · 10.0s turn · 10.0s LLM · 6,120 chars (click to expand)
turn 310.0s turn · 10.0s LLM · 0.0s tools
call → run_correctness
callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out,\n int rows, int scalar_in, int scalar_out_dim, int vector_in, int vector_out_dim)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n int C = scalar_in + vector_in;\n extern __shared__ float sh[];\n float* comb = sh;\n float* gate = sh + C;\n float* vec = gate + vector_out_dim;\n\n const float* srow = scalar_feat + (size_t)row * scalar_in;\n const float* vrow = vector_feat + (size_t)row * vector_in * 3;\n\n for (int i = tid; i < scalar_in; i += blockDim.x) comb[i] = srow[i];\n for (int i = tid; i < vector_in * 3; i += blockDim.x) vec[i] = vrow[i];\n __syncthreads();\n\n for (int i = tid; i < vector_in; i += blockDim.x) {\n float x = vec[i*3 + 0], y = vec[i*3 + 1], z = vec[i*3 + 2];\n comb[scalar_in + i] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n for (int o = tid; o < scalar_out_dim; o += blockDim.x) {\n const float* w = scalar_w + (size_t)o * C;\n float acc = scalar_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) acc = fmaf(w[i], comb[i], acc);\n scalar_out[(size_t)row * scalar_out_dim + o] = acc;\n }\n\n for (int o = tid; o < vector_out_dim; o += blockDim.x) {\n const float* w = gate_w + (size_t)o * C;\n float acc = gate_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) acc = fmaf(w[i], comb[i], acc);\n gate[o] = 1.0f / (1.0f + expf(-acc));\n }\n __syncthreads();\n\n int total_v = vector_out_dim * 3;\n for (int idx = tid; idx < total_v; idx += blockDim.x) {\n int o = idx / 3;\n int d = idx - o * 3;\n const float* w = vector_w + (size_t)o * vector_in;\n float acc = 0.0f;\n #pragma unroll 4\n for (int i = 0; i < vector_in; ++i) acc = fmaf(w[i], vec[i*3 + d], acc);\n vector_out[((size_t)row * vector_out_dim + o) * 3 + d] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n int scalar_in = (int)scalar_feat.size(2);\n int vector_in = (int)vector_feat.size(2);\n int scalar_out_dim = (int)scalar_w.size(0);\n int vector_out_dim = (int)vector_w.size(0);\n auto s_out = torch::empty({B, N, scalar_out_dim}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_out_dim, 3}, vector_feat.options());\n int rows = (int)(B * N);\n int threads = 128;\n size_t shmem = (size_t)(scalar_in + vector_in + vector_out_dim + vector_in * 3) * sizeof(float);\n se3_forward_kernel<<<rows, threads, shmem, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(),\n vector_w.data_ptr<float>(), gate_w.data_ptr<float>(), gate_b.data_ptr<float>(),\n s_out.data_ptr<float>(), v_out.data_ptr<float>(),\n rows, scalar_in, scalar_out_dim, vector_in, vector_out_dim);\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(\n name='se3_fused_ext_v2',\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=['se3_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat,\n self.scalar_linear.weight, self.scalar_linear.bias,\n self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=7.15e-07 mean_abs_err=4.15e-08 max_rel_err=5.28e-02
turn 4 · 1 tool call(s) · 15.4s turn · 15.4s LLM · 7,188 chars (click to expand)
turn 415.4s turn · 15.4s LLM · 0.0s tools
reasoning
**Optimizing kernel performance** I need to figure out how to optimize the kernel for better performance, considering the timing. Currently, we don’t have an effective timing tool, so I only have the submit option. I’m thinking about how to improve the kernel design, possibly by using fixed dimensions to avoid dynamic loops with specialized unrolling. I could assume known dimensions and use the host to call specialized kernels, which might speed things up. Then, I’ll also consider warp-level matrix multiplication for efficiency. **Analyzing H100 performance** It looks like on the H100, I might be able to achieve around 1,024 blocks with approximately 6 million operations, possibly launching them in about 10 microseconds. The reference could be over 100 microseconds due to operations, indicating there's room for speedup. I wonder if I can improve performance by combining two rows per block. While 1,024 blocks seems fine, each block has 128 threads with only 80 active for the main dot operation. I need to ensure I'm correctly using fast math for accuracy checks.
call → static_check
callstatic_check { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out,\n int rows, int scalar_in, int scalar_out_dim, int vector_in, int vector_out_dim)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n int C = scalar_in + vector_in;\n extern __shared__ float sh[];\n float* comb = sh;\n float* gate = sh + C;\n float* vec = gate + vector_out_dim;\n\n const float* srow = scalar_feat + (size_t)row * scalar_in;\n const float* vrow = vector_feat + (size_t)row * vector_in * 3;\n\n for (int i = tid; i < scalar_in; i += blockDim.x) comb[i] = srow[i];\n for (int i = tid; i < vector_in * 3; i += blockDim.x) vec[i] = vrow[i];\n __syncthreads();\n\n for (int i = tid; i < vector_in; i += blockDim.x) {\n float x = vec[i*3 + 0], y = vec[i*3 + 1], z = vec[i*3 + 2];\n comb[scalar_in + i] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n for (int o = tid; o < scalar_out_dim; o += blockDim.x) {\n const float* w = scalar_w + (size_t)o * C;\n float acc = scalar_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) acc = fmaf(w[i], comb[i], acc);\n scalar_out[(size_t)row * scalar_out_dim + o] = acc;\n }\n\n for (int o = tid; o < vector_out_dim; o += blockDim.x) {\n const float* w = gate_w + (size_t)o * C;\n float acc = gate_b[o];\n #pragma unroll 4\n for (int i = 0; i < C; ++i) acc = fmaf(w[i], comb[i], acc);\n gate[o] = 1.0f / (1.0f + expf(-acc));\n }\n __syncthreads();\n\n int total_v = vector_out_dim * 3;\n for (int idx = tid; idx < total_v; idx += blockDim.x) {\n int o = idx / 3;\n int d = idx - o * 3;\n const float* w = vector_w + (size_t)o * vector_in;\n float acc = 0.0f;\n #pragma unroll 4\n for (int i = 0; i < vector_in; ++i) acc = fmaf(w[i], vec[i*3 + d], acc);\n vector_out[((size_t)row * vector_out_dim + o) * 3 + d] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat,\n torch::Tensor vector_feat,\n torch::Tensor scalar_w,\n torch::Tensor scalar_b,\n torch::Tensor vector_w,\n torch::Tensor gate_w,\n torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n int scalar_in = (int)scalar_feat.size(2);\n int vector_in = (int)vector_feat.size(2);\n int scalar_out_dim = (int)scalar_w.size(0);\n int vector_out_dim = (int)vector_w.size(0);\n auto s_out = torch::empty({B, N, scalar_out_dim}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_out_dim, 3}, vector_feat.options());\n int rows = (int)(B * N);\n int threads = 128;\n size_t shmem = (size_t)(scalar_in + vector_in + vector_out_dim + vector_in * 3) * sizeof(float);\n se3_forward_kernel<<<rows, threads, shmem, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(),\n vector_w.data_ptr<float>(), gate_w.data_ptr<float>(), gate_b.data_ptr<float>(),\n s_out.data_ptr<float>(), v_out.data_ptr<float>(),\n rows, scalar_in, scalar_out_dim, vector_in, vector_out_dim);\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(\n name='se3_fused_ext_v2',\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=['se3_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat,\n self.scalar_linear.weight, self.scalar_linear.bias,\n self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← static_check (ok)
okstatic_check static_check PASSED: no violations or warnings detected.
turn 513.6s turn · 13.6s LLM · 0.0s tools
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_64_16_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n __shared__ float comb[80];\n __shared__ float gate[16];\n __shared__ float vec[48];\n\n const float* srow = scalar_feat + ((size_t)row << 6); // *64\n const float* vrow = vector_feat + (size_t)row * 48;\n\n if (tid < 64) comb[tid] = srow[tid];\n if (tid < 48) vec[tid] = vrow[tid];\n __syncthreads();\n\n if (tid < 16) {\n float x = vec[tid*3 + 0];\n float y = vec[tid*3 + 1];\n float z = vec[tid*3 + 2];\n comb[64 + tid] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n if (tid < 80) {\n int o = tid;\n float acc;\n const float* w;\n if (o < 64) {\n w = scalar_w + (size_t)o * 80;\n acc = scalar_b[o];\n } else {\n int g = o - 64;\n w = gate_w + (size_t)g * 80;\n acc = gate_b[g];\n }\n #pragma unroll\n for (int i = 0; i < 80; ++i) {\n acc = fmaf(w[i], comb[i], acc);\n }\n if (o < 64) {\n scalar_out[((size_t)row << 6) + o] = acc;\n } else {\n int g = o - 64;\n gate[g] = 1.0f / (1.0f + expf(-acc));\n }\n }\n __syncthreads();\n\n if (tid < 48) {\n int o = tid / 3;\n int d = tid - o * 3;\n const float* w = vector_w + ((size_t)o << 4); // *16\n float acc = 0.0f;\n #pragma unroll\n for (int i = 0; i < 16; ++i) {\n acc = fmaf(w[i], vec[i*3 + d], acc);\n }\n vector_out[(size_t)row * 48 + tid] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n auto s_out = torch::empty({B, N, scalar_w.size(0)}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_w.size(0), 3}, vector_feat.options());\n int rows = (int)(B * N);\n se3_forward_64_16_kernel<<<rows, 128, 0, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(), vector_w.data_ptr<float>(),\n gate_w.data_ptr<float>(), gate_b.data_ptr<float>(), s_out.data_ptr<float>(), v_out.data_ptr<float>());\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(name='se3_fused_ext_v3', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['se3_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat, self.scalar_linear.weight,\n self.scalar_linear.bias, self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 69.3s turn · 9.3s LLM · 0.0s tools
call → run_correctness
callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_64_16_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n __shared__ float comb[80];\n __shared__ float gate[16];\n __shared__ float vec[48];\n\n const float* srow = scalar_feat + ((size_t)row << 6);\n const float* vrow = vector_feat + (size_t)row * 48;\n\n if (tid < 64) comb[tid] = srow[tid];\n if (tid < 48) vec[tid] = vrow[tid];\n __syncthreads();\n\n if (tid < 16) {\n float x = vec[tid*3 + 0];\n float y = vec[tid*3 + 1];\n float z = vec[tid*3 + 2];\n comb[64 + tid] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n if (tid < 80) {\n int o = tid;\n float acc;\n const float* w;\n if (o < 64) {\n w = scalar_w + (size_t)o * 80;\n acc = scalar_b[o];\n } else {\n int g = o - 64;\n w = gate_w + (size_t)g * 80;\n acc = gate_b[g];\n }\n #pragma unroll\n for (int i = 0; i < 80; ++i) {\n acc = fmaf(w[i], comb[i], acc);\n }\n if (o < 64) {\n scalar_out[((size_t)row << 6) + o] = acc;\n } else {\n int g = o - 64;\n gate[g] = 1.0f / (1.0f + expf(-acc));\n }\n }\n __syncthreads();\n\n if (tid < 48) {\n int o = tid / 3;\n int d = tid - o * 3;\n const float* w = vector_w + ((size_t)o << 4);\n float acc = 0.0f;\n #pragma unroll\n for (int i = 0; i < 16; ++i) {\n acc = fmaf(w[i], vec[i*3 + d], acc);\n }\n vector_out[(size_t)row * 48 + tid] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n auto s_out = torch::empty({B, N, scalar_w.size(0)}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_w.size(0), 3}, vector_feat.options());\n int rows = (int)(B * N);\n se3_forward_64_16_kernel<<<rows, 128, 0, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(), vector_w.data_ptr<float>(),\n gate_w.data_ptr<float>(), gate_b.data_ptr<float>(), s_out.data_ptr<float>(), v_out.data_ptr<float>());\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(name='se3_fused_ext_v3', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['se3_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat, self.scalar_linear.weight,\n self.scalar_linear.bias, self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=7.15e-07 mean_abs_err=4.15e-08 max_rel_err=5.28e-02
turn 715.1s turn · 15.1s LLM · 0.0s tools
call → static_check
callstatic_check { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_64_16_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n __shared__ float comb[80];\n __shared__ float gate[16];\n __shared__ float vec[48];\n\n const float* srow = scalar_feat + ((size_t)row << 6);\n const float* vrow = vector_feat + (size_t)row * 48;\n\n if (tid < 64) comb[tid] = srow[tid];\n if (tid < 48) vec[tid] = vrow[tid];\n __syncthreads();\n\n if (tid < 16) {\n float x = vec[tid*3 + 0];\n float y = vec[tid*3 + 1];\n float z = vec[tid*3 + 2];\n comb[64 + tid] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n if (tid < 80) {\n int o = tid;\n float acc;\n const float* w;\n if (o < 64) {\n w = scalar_w + (size_t)o * 80;\n acc = scalar_b[o];\n } else {\n int g = o - 64;\n w = gate_w + (size_t)g * 80;\n acc = gate_b[g];\n }\n #pragma unroll\n for (int i = 0; i < 80; ++i) {\n acc = fmaf(w[i], comb[i], acc);\n }\n if (o < 64) {\n scalar_out[((size_t)row << 6) + o] = acc;\n } else {\n int g = o - 64;\n gate[g] = 1.0f / (1.0f + expf(-acc));\n }\n }\n __syncthreads();\n\n if (tid < 48) {\n int o = tid / 3;\n int d = tid - o * 3;\n const float* w = vector_w + ((size_t)o << 4);\n float acc = 0.0f;\n #pragma unroll\n for (int i = 0; i < 16; ++i) {\n acc = fmaf(w[i], vec[i*3 + d], acc);\n }\n vector_out[(size_t)row * 48 + tid] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n auto s_out = torch::empty({B, N, scalar_w.size(0)}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_w.size(0), 3}, vector_feat.options());\n int rows = (int)(B * N);\n se3_forward_64_16_kernel<<<rows, 128, 0, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(), vector_w.data_ptr<float>(),\n gate_w.data_ptr<float>(), gate_b.data_ptr<float>(), s_out.data_ptr<float>(), v_out.data_ptr<float>());\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(name='se3_fused_ext_v3', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['se3_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat, self.scalar_linear.weight,\n self.scalar_linear.bias, self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← static_check (ok)
okstatic_check static_check PASSED: no violations or warnings detected.
turn 8 · FINAL · 1 tool call(s) · 8.7s turn · 8.7s LLM · 6,014 chars (click to expand)
turn 88.7s turn · 8.7s LLM · 0.0s tools · FINAL
reasoning
**Finalizing submission** I’m preparing to submit, and it looks like I should finalize version 3. I need to make sure to include the code that’s essential for it. I think I need to double-check everything before hitting submit to avoid any mistakes. Clarity is key here, and I want to ensure all necessary components are present. Let's make sure it’s all ready to go!
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n#include <vector>\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__global__ void se3_forward_64_16_kernel(\n const float* __restrict__ scalar_feat,\n const float* __restrict__ vector_feat,\n const float* __restrict__ scalar_w,\n const float* __restrict__ scalar_b,\n const float* __restrict__ vector_w,\n const float* __restrict__ gate_w,\n const float* __restrict__ gate_b,\n float* __restrict__ scalar_out,\n float* __restrict__ vector_out)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n __shared__ float comb[80];\n __shared__ float gate[16];\n __shared__ float vec[48];\n\n const float* srow = scalar_feat + ((size_t)row << 6);\n const float* vrow = vector_feat + (size_t)row * 48;\n\n if (tid < 64) comb[tid] = srow[tid];\n if (tid < 48) vec[tid] = vrow[tid];\n __syncthreads();\n\n if (tid < 16) {\n float x = vec[tid*3 + 0];\n float y = vec[tid*3 + 1];\n float z = vec[tid*3 + 2];\n comb[64 + tid] = sqrtf(x*x + y*y + z*z);\n }\n __syncthreads();\n\n if (tid < 80) {\n int o = tid;\n float acc;\n const float* w;\n if (o < 64) {\n w = scalar_w + (size_t)o * 80;\n acc = scalar_b[o];\n } else {\n int g = o - 64;\n w = gate_w + (size_t)g * 80;\n acc = gate_b[g];\n }\n #pragma unroll\n for (int i = 0; i < 80; ++i) {\n acc = fmaf(w[i], comb[i], acc);\n }\n if (o < 64) {\n scalar_out[((size_t)row << 6) + o] = acc;\n } else {\n int g = o - 64;\n gate[g] = 1.0f / (1.0f + expf(-acc));\n }\n }\n __syncthreads();\n\n if (tid < 48) {\n int o = tid / 3;\n int d = tid - o * 3;\n const float* w = vector_w + ((size_t)o << 4);\n float acc = 0.0f;\n #pragma unroll\n for (int i = 0; i < 16; ++i) {\n acc = fmaf(w[i], vec[i*3 + d], acc);\n }\n vector_out[(size_t)row * 48 + tid] = gate[o] * acc;\n }\n}\n\nstd::vector<torch::Tensor> se3_forward(\n torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,\n torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b)\n{\n int64_t B = scalar_feat.size(0);\n int64_t N = scalar_feat.size(1);\n auto s_out = torch::empty({B, N, scalar_w.size(0)}, scalar_feat.options());\n auto v_out = torch::empty({B, N, vector_w.size(0), 3}, vector_feat.options());\n int rows = (int)(B * N);\n se3_forward_64_16_kernel<<<rows, 128, 0, at::cuda::getCurrentCUDAStream()>>>(\n scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),\n scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(), vector_w.data_ptr<float>(),\n gate_w.data_ptr<float>(), gate_b.data_ptr<float>(), s_out.data_ptr<float>(), v_out.data_ptr<float>());\n return {s_out, v_out};\n}\n'''\n\n_ext = load_inline(name='se3_fused_ext_v3', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['se3_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, scalar_in, scalar_out, vector_in, vector_out):\n super().__init__()\n self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)\n self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))\n nn.init.xavier_uniform_(self.vector_weight)\n self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)\n def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:\n return tuple(_ext.se3_forward(scalar_feat, vector_feat, self.scalar_linear.weight,\n self.scalar_linear.bias, self.vector_weight,\n self.vector_gate.weight, self.vector_gate.bias))\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.02 μs Runtime stats: mean=0.02μs median=?μs std=0.01μs Numerical precision: max_abs_err=7.15e-07 mean_abs_err=4.15e-08 max_rel_err=5.28e-02 Memory: 33.93 MB peak (ref: 34.55 MB, ratio: 0.9819x) Kernel launches: 2 (ref: 51, fusion ratio: 25.5) SOL score: 0.0079
final submitted kernel (122 lines, 4,313 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_src = r'''
#include <torch/extension.h>
#include <vector>
std::vector<torch::Tensor> se3_forward(
    torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,
    torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b);
'''

_cuda_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <vector>

__global__ void se3_forward_64_16_kernel(
    const float* __restrict__ scalar_feat,
    const float* __restrict__ vector_feat,
    const float* __restrict__ scalar_w,
    const float* __restrict__ scalar_b,
    const float* __restrict__ vector_w,
    const float* __restrict__ gate_w,
    const float* __restrict__ gate_b,
    float* __restrict__ scalar_out,
    float* __restrict__ vector_out)
{
    int row = blockIdx.x;
    int tid = threadIdx.x;
    __shared__ float comb[80];
    __shared__ float gate[16];
    __shared__ float vec[48];

    const float* srow = scalar_feat + ((size_t)row << 6);
    const float* vrow = vector_feat + (size_t)row * 48;

    if (tid < 64) comb[tid] = srow[tid];
    if (tid < 48) vec[tid] = vrow[tid];
    __syncthreads();

    if (tid < 16) {
        float x = vec[tid*3 + 0];
        float y = vec[tid*3 + 1];
        float z = vec[tid*3 + 2];
        comb[64 + tid] = sqrtf(x*x + y*y + z*z);
    }
    __syncthreads();

    if (tid < 80) {
        int o = tid;
        float acc;
        const float* w;
        if (o < 64) {
            w = scalar_w + (size_t)o * 80;
            acc = scalar_b[o];
        } else {
            int g = o - 64;
            w = gate_w + (size_t)g * 80;
            acc = gate_b[g];
        }
        #pragma unroll
        for (int i = 0; i < 80; ++i) {
            acc = fmaf(w[i], comb[i], acc);
        }
        if (o < 64) {
            scalar_out[((size_t)row << 6) + o] = acc;
        } else {
            int g = o - 64;
            gate[g] = 1.0f / (1.0f + expf(-acc));
        }
    }
    __syncthreads();

    if (tid < 48) {
        int o = tid / 3;
        int d = tid - o * 3;
        const float* w = vector_w + ((size_t)o << 4);
        float acc = 0.0f;
        #pragma unroll
        for (int i = 0; i < 16; ++i) {
            acc = fmaf(w[i], vec[i*3 + d], acc);
        }
        vector_out[(size_t)row * 48 + tid] = gate[o] * acc;
    }
}

std::vector<torch::Tensor> se3_forward(
    torch::Tensor scalar_feat, torch::Tensor vector_feat, torch::Tensor scalar_w,
    torch::Tensor scalar_b, torch::Tensor vector_w, torch::Tensor gate_w, torch::Tensor gate_b)
{
    int64_t B = scalar_feat.size(0);
    int64_t N = scalar_feat.size(1);
    auto s_out = torch::empty({B, N, scalar_w.size(0)}, scalar_feat.options());
    auto v_out = torch::empty({B, N, vector_w.size(0), 3}, vector_feat.options());
    int rows = (int)(B * N);
    se3_forward_64_16_kernel<<<rows, 128, 0, at::cuda::getCurrentCUDAStream()>>>(
        scalar_feat.data_ptr<float>(), vector_feat.data_ptr<float>(),
        scalar_w.data_ptr<float>(), scalar_b.data_ptr<float>(), vector_w.data_ptr<float>(),
        gate_w.data_ptr<float>(), gate_b.data_ptr<float>(), s_out.data_ptr<float>(), v_out.data_ptr<float>());
    return {s_out, v_out};
}
'''

_ext = load_inline(name='se3_fused_ext_v3', cpp_sources=_cpp_src, cuda_sources=_cuda_src,
                   functions=['se3_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'],
                   extra_cflags=['-O3'], verbose=False)

class ModelNew(nn.Module):
    def __init__(self, scalar_in, scalar_out, vector_in, vector_out):
        super().__init__()
        self.scalar_linear = nn.Linear(scalar_in + vector_in, scalar_out)
        self.vector_weight = nn.Parameter(torch.randn(vector_out, vector_in))
        nn.init.xavier_uniform_(self.vector_weight)
        self.vector_gate = nn.Linear(scalar_in + vector_in, vector_out)
    def forward(self, scalar_feat: torch.Tensor, vector_feat: torch.Tensor) -> tuple:
        return tuple(_ext.se3_forward(scalar_feat, vector_feat, self.scalar_linear.weight,
                                      self.scalar_linear.bias, self.vector_weight,
                                      self.vector_gate.weight, self.vector_gate.bias))