correct
22_SE3InvariantLinear.py
started
2026-05-02T17:49:15Z
timingkernelbench.timing
kernel mean / median / std
0.02 / — / 0.01 μs
ref mean / median / std
0.10 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
25.50x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
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))