gpt-5.5-priority · popcorn · l123_single_gpt
correct 36_InvariantPointAttention.py
model
gpt-5.5-priority
turns / max
1 / 1
tool calls
1
tokens (total)
0
wall clock
—
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.05 μs
ref runtime
0.80 μs
speedup
16.53x
started
2026-05-05T00:01:59Z
timingkernelbench.timing
kernel runtime
0.05 μs
ref runtime
0.80 μs
speedup
16.53x
kernel mean / median / std
0.05 / — / 5.21e-04 μs
ref mean / median / std
0.80 / — / 0.04 μs
numerical precisioneval_kernel_against_ref
max abs error
5.07e-07
mean abs error
8.10e-08
max rel error
0.01
mean rel error
2.57e-06
gpu memoryextended_metrics.measure_memory
kernel peak
33.69 MB
ref peak
34.48 MB
kernel / ref
0.98x
kernel launches / fusiontorch.profiler
kernel launches
8
ref launches
309
fusion ratio (ref / kernel)
38.62x
total cuda time
64 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.011
dram util %
1.1%
compute util %
0.0%
bottleneck
unknown
dominant pipe
None
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
38.10 GB/s
dram util %
1.1%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
14.0%
bottleneck
unknown
dominant pipe %
—
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
286.4%
ld sectors / req
52.36
st sectors / req
—
regs / thread
164
smem / block
4608 B
block size
256
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
6005.7%
short_scoreboard
156.4%
flagssubmit_kernel guard
excessive speedup
flagged
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F
import math


class Model(nn.Module):
    """
    Invariant Point Attention (IPA) from AlphaFold2's structure module.
    Combines sequence-space attention with 3-D geometric queries/keys/values
    expressed in local residue frames.  The attention logits incorporate
    Euclidean distances in global frame, making it SE(3)-equivariant.
    """

    def __init__(self, node_dim, pair_dim, num_heads, num_query_points, num_value_points):
        super().__init__()
        self.num_heads = num_heads
        self.head_dim = node_dim // num_heads
        self.num_query_points = num_query_points
        self.num_value_points = num_value_points
        assert node_dim % num_heads == 0

        d = self.head_dim
        self.q_proj = nn.Linear(node_dim, num_heads * d, bias=False)
        self.k_proj = nn.Linear(node_dim, num_heads * d, bias=False)
        self.v_proj = nn.Linear(node_dim, num_heads * d, bias=False)

        self.q_pts = nn.Linear(node_dim, num_heads * num_query_points * 3, bias=False)
        self.k_pts = nn.Linear(node_dim, num_heads * num_query_points * 3, bias=False)
        self.v_pts = nn.Linear(node_dim, num_heads * num_value_points * 3, bias=False)

        self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)

        self.head_weights = nn.Parameter(torch.zeros(num_heads))

        out_features = num_heads * d + num_heads * num_value_points * 3
        self.out_proj = nn.Linear(out_features, node_dim)

        self.layer_norm = nn.LayerNorm(node_dim)

    def _apply_frame(self, pts, rotations, translations):
        """Transform local-frame points to global frame."""
        # pts: (B, N, P, 3), rotations: (B, N, 3, 3), translations: (B, N, 3)
        global_pts = torch.einsum("bnpc,bncd->bnpd", pts, rotations) + translations.unsqueeze(2)
        return global_pts

    def _apply_inv_frame(self, pts, rotations, translations):
        """Transform global-frame points to local frame."""
        centered = pts - translations.unsqueeze(2)
        local_pts = torch.einsum("bnpc,bndc->bnpd", centered, rotations)
        return local_pts

    def forward(
        self,
        node_repr: torch.Tensor,
        pair_repr: torch.Tensor,
        translations: torch.Tensor,
        rotations: torch.Tensor,
    ) -> torch.Tensor:
        """
        Args:
            node_repr:    (B, N, node_dim)
            pair_repr:    (B, N, N, pair_dim)
            translations: (B, N, 3)
            rotations:    (B, N, 3, 3)
        Returns:
            updated node representation (B, N, node_dim)
        """
        B, N, D = node_repr.shape
        h = self.num_heads
        d = self.head_dim
        Qp = self.num_query_points
        Vp = self.num_value_points

        x = self.layer_norm(node_repr)

        q = self.q_proj(x).view(B, N, h, d)
        k = self.k_proj(x).view(B, N, h, d)
        v = self.v_proj(x).view(B, N, h, d)

        # Points in local frame, then transform to global
        q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)
        k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)
        v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)

        q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)
        k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)
        v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)

        # Scalar attention logits
        attn_scalar = torch.einsum("bihd,bjhd->bhij", q, k) / math.sqrt(d)

        # Point attention: negative squared distance between query and key points
        # q_pts: (B, N_i, h, Qp, 3), k_pts: (B, N_j, h, Qp, 3)
        q_expand = q_pts.unsqueeze(2)  # (B, N, 1, h, Qp, 3)
        k_expand = k_pts.unsqueeze(1)  # (B, 1, N, h, Qp, 3)
        pt_dist = ((q_expand - k_expand) ** 2).sum(dim=(-1, -2))  # (B, N, N, h)

        w = F.softplus(self.head_weights)
        pt_attn = -(w.view(1, 1, 1, h) * pt_dist) / 2.0
        pt_attn = pt_attn.permute(0, 3, 1, 2)  # (B, h, N, N)

        # Pair bias
        pair_bias = self.pair_bias(pair_repr).permute(0, 3, 1, 2)  # (B, h, N, N)

        attn = attn_scalar + pt_attn + pair_bias
        attn = F.softmax(attn, dim=-1)  # (B, h, N_i, N_j)

        # Aggregate scalar values
        out_scalar = torch.einsum("bhij,bjhd->bihd", attn, v).reshape(B, N, h * d)

        # Aggregate point values in global frame, then transform back to local
        out_pts = torch.einsum("bhij,bjhpc->bihpc", attn, v_pts)  # (B, N, h, Vp, 3)
        out_pts = out_pts.reshape(B, N, h * Vp, 3)
        out_pts_local = self._apply_inv_frame(out_pts, rotations, translations)
        out_pts_flat = out_pts_local.reshape(B, N, h * Vp * 3)

        out = torch.cat([out_scalar, out_pts_flat], dim=-1)
        return self.out_proj(out)


node_dim = 64
pair_dim = 32
num_heads = 4
num_query_points = 4
num_value_points = 8
seq_len = 32
batch_size = 2


def get_inputs():
    p = popcorn_pri
    node_repr = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), node_dim)
    pair_repr = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), p.jitter_int(seq_len), pair_dim)
    translations = torch.randn(p.jitter_int(batch_size), p.jitter_int(seq_len), 3)
    rotations = torch.zeros(p.jitter_int(batch_size), p.jitter_int(seq_len), 3, 3)
    for b in range(p.jitter_int(batch_size)):
        for n in range(p.jitter_int(seq_len)):
            rotations[b, n] = torch.eye(3)
    return [node_repr, pair_repr, translations, rotations]


def get_init_inputs():
    return [node_dim, pair_dim, num_heads, num_query_points, num_value_points]
turn 0 · FINAL · 1 tool call(s) · 131.7s turn · 131.7s LLM · 15,947 chars (click to expand)
turn 0131.7s turn · 131.7s LLM · 0.0s tools · FINAL
reasoning
**Planning a Python model** I need to craft a Python file and compile it inline. I'm thinking about fixed dimensions for our model. The constructor arguments are crucial; should I use weights from nn.Linear? ModelNew should match parameters for consistent initialization. Evaluating might copy weights, so defining identical layers is important for the state_dict. The input dimensions are fixed, but various arguments are accepted. I'm also considering using torch operations for projections and possibly implementing a custom CUDA kernel for the forward process to maximize speed. **Designing a CUDA model** I'm working with small sizes: nodes at 64 and calculations for B*N yield 64 rows, which suggests small gemm operations. There’s some overhead from PyTorch, so I’m considering a custom fused implementation. I need to create a CUDA extension for the function ipa_forward, where the host allocates the output. I'll implement layer normalization and projections within a fused kernel per batch size (B, N). Weights will be organized according to nn.Linear standards, and I must handle output computations efficiently. The attention output needs careful aggregation too! **Streamlining parameters** I need to set up scalar values for q and k, plus global values for q_pts and k_pts, as well as the scalar for v and global v_pts. It seems efficient to fuse all of these computations into a single block for each (B, i, h). This approach should optimize performance while keeping the operations organized. I have to ensure that everything is perfectly aligned for maximum efficiency in the calculations.
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\ntorch::Tensor ipa_forward(\n torch::Tensor node_repr, torch::Tensor pair_repr,\n torch::Tensor translations, torch::Tensor rotations,\n torch::Tensor ln_w, torch::Tensor ln_b,\n torch::Tensor q_w, torch::Tensor k_w, torch::Tensor v_w,\n torch::Tensor qpts_w, torch::Tensor kpts_w, torch::Tensor vpts_w,\n torch::Tensor pair_w, torch::Tensor head_w,\n torch::Tensor out_w, torch::Tensor out_b,\n int64_t num_heads, int64_t num_query_points, int64_t num_value_points);\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 <math.h>\n\nstatic __device__ __forceinline__ float softplus_dev(float x) {\n if (x > 20.0f) return x;\n if (x < -20.0f) return expf(x);\n return log1pf(expf(x));\n}\n\n__global__ void ln_proj_kernel(\n const float* __restrict__ node,\n const float* __restrict__ ln_w,\n const float* __restrict__ ln_b,\n const float* __restrict__ q_w,\n const float* __restrict__ k_w,\n const float* __restrict__ v_w,\n const float* __restrict__ qpts_w,\n const float* __restrict__ kpts_w,\n const float* __restrict__ vpts_w,\n float* __restrict__ q,\n float* __restrict__ k,\n float* __restrict__ v,\n float* __restrict__ qpl,\n float* __restrict__ kpl,\n float* __restrict__ vpl,\n int BN, int D, int H, int QP, int VP)\n{\n int row = blockIdx.x;\n int tid = threadIdx.x;\n __shared__ float red[128];\n __shared__ float xn[128];\n\n float s = 0.0f;\n for (int i = tid; i < D; i += blockDim.x) s += node[row * D + i];\n red[tid] = s;\n __syncthreads();\n for (int stride = blockDim.x/2; stride > 0; stride >>= 1) {\n if (tid < stride) red[tid] += red[tid + stride];\n __syncthreads();\n }\n float mean = red[0] / (float)D;\n\n s = 0.0f;\n for (int i = tid; i < D; i += blockDim.x) {\n float a = node[row * D + i] - mean;\n s += a * a;\n }\n red[tid] = s;\n __syncthreads();\n for (int stride = blockDim.x/2; stride > 0; stride >>= 1) {\n if (tid < stride) red[tid] += red[tid + stride];\n __syncthreads();\n }\n float inv_std = rsqrtf(red[0] / (float)D + 1.0e-5f);\n for (int i = tid; i < D; i += blockDim.x) {\n xn[i] = (node[row * D + i] - mean) * inv_std * ln_w[i] + ln_b[i];\n }\n __syncthreads();\n\n int dhead = D / H;\n int qkv = H * dhead;\n int qpo = H * QP * 3;\n int vpo = H * VP * 3;\n int total = 3*qkv + 2*qpo + vpo;\n\n for (int o = tid; o < total; o += blockDim.x) {\n const float* W;\n float* O;\n int oo, out_stride;\n if (o < qkv) { W = q_w; O = q; oo = o; out_stride = qkv; }\n else if (o < 2*qkv) { W = k_w; O = k; oo = o - qkv; out_stride = qkv; }\n else if (o < 3*qkv) { W = v_w; O = v; oo = o - 2*qkv; out_stride = qkv; }\n else if (o < 3*qkv + qpo) { W = qpts_w; O = qpl; oo = o - 3*qkv; out_stride = qpo; }\n else if (o < 3*qkv + 2*qpo) { W = kpts_w; O = kpl; oo = o - 3*qkv - qpo; out_stride = qpo; }\n else { W = vpts_w; O = vpl; oo = o - 3*qkv - 2*qpo; out_stride = vpo; }\n float acc = 0.0f;\n const float* wrow = W + oo * D;\n #pragma unroll\n for (int ii = 0; ii < 64; ++ii) {\n if (ii < D) acc += xn[ii] * wrow[ii];\n }\n O[row * out_stride + oo] = acc;\n }\n}\n\n__global__ void attn_kernel(\n const float* __restrict__ pair_repr,\n const float* __restrict__ translations,\n const float* __restrict__ rotations,\n const float* __restrict__ q,\n const float* __restrict__ k,\n const float* __restrict__ v,\n const float* __restrict__ qpl,\n const float* __restrict__ kpl,\n const float* __restrict__ vpl,\n const float* __restrict__ pair_w,\n const float* __restrict__ head_w,\n float* __restrict__ cat,\n int B, int N, int D, int PD, int H, int QP, int VP)\n{\n int bid = blockIdx.x;\n int h = bid % H;\n int i = (bid / H) % N;\n int b = bid / (H * N);\n int tid = threadIdx.x;\n int dhead = D / H;\n int qkv_stride = H * dhead;\n int qp_stride = H * QP * 3;\n int vp_stride = H * VP * 3;\n int row_i = b * N + i;\n\n __shared__ float logits[64];\n __shared__ float attn[64];\n\n if (tid < N) {\n int j = tid;\n int row_j = b * N + j;\n float sc = 0.0f;\n for (int dd = 0; dd < dhead; ++dd) {\n sc += q[row_i * qkv_stride + h * dhead + dd] * k[row_j * qkv_stride + h * dhead + dd];\n }\n sc *= rsqrtf((float)dhead);\n\n float dist = 0.0f;\n for (int p = 0; p < QP; ++p) {\n float ql0 = qpl[row_i * qp_stride + (h*QP + p)*3 + 0];\n float ql1 = qpl[row_i * qp_stride + (h*QP + p)*3 + 1];\n float ql2 = qpl[row_i * qp_stride + (h*QP + p)*3 + 2];\n float kl0 = kpl[row_j * qp_stride + (h*QP + p)*3 + 0];\n float kl1 = kpl[row_j * qp_stride + (h*QP + p)*3 + 1];\n float kl2 = kpl[row_j * qp_stride + (h*QP + p)*3 + 2];\n for (int c = 0; c < 3; ++c) {\n float qg = ql0 * rotations[(row_i*3 + 0)*3 + c] + ql1 * rotations[(row_i*3 + 1)*3 + c] + ql2 * rotations[(row_i*3 + 2)*3 + c] + translations[row_i*3 + c];\n float kg = kl0 * rotations[(row_j*3 + 0)*3 + c] + kl1 * rotations[(row_j*3 + 1)*3 + c] + kl2 * rotations[(row_j*3 + 2)*3 + c] + translations[row_j*3 + c];\n float diff = qg - kg;\n dist += diff * diff;\n }\n }\n float pb = 0.0f;\n const float* pr = pair_repr + ((b * N + i) * N + j) * PD;\n const float* pw = pair_w + h * PD;\n for (int pp = 0; pp < PD; ++pp) pb += pr[pp] * pw[pp];\n float w = softplus_dev(head_w[h]);\n logits[j] = sc + pb - 0.5f * w * dist;\n }\n __syncthreads();\n\n if (tid == 0) {\n float m = logits[0];\n for (int j = 1; j < N; ++j) m = fmaxf(m, logits[j]);\n float s = 0.0f;\n for (int j = 0; j < N; ++j) {\n float e = expf(logits[j] - m);\n attn[j] = e;\n s += e;\n }\n float invs = 1.0f / s;\n for (int j = 0; j < N; ++j) attn[j] *= invs;\n }\n __syncthreads();\n\n int out_width = H * dhead + H * VP * 3;\n if (tid < dhead) {\n float acc = 0.0f;\n for (int j = 0; j < N; ++j) {\n int row_j = b * N + j;\n acc += attn[j] * v[row_j * qkv_stride + h * dhead + tid];\n }\n cat[row_i * out_width + h * dhead + tid] = acc;\n }\n if (tid < VP * 3) {\n int vp = tid / 3;\n int ld = tid - vp * 3;\n float glob[3];\n for (int c = 0; c < 3; ++c) {\n float acc = 0.0f;\n for (int j = 0; j < N; ++j) {\n int row_j = b * N + j;\n float vl0 = vpl[row_j * vp_stride + (h*VP + vp)*3 + 0];\n float vl1 = vpl[row_j * vp_stride + (h*VP + vp)*3 + 1];\n float vl2 = vpl[row_j * vp_stride + (h*VP + vp)*3 + 2];\n float vg = vl0 * rotations[(row_j*3 + 0)*3 + c] + vl1 * rotations[(row_j*3 + 1)*3 + c] + vl2 * rotations[(row_j*3 + 2)*3 + c] + translations[row_j*3 + c];\n acc += attn[j] * vg;\n }\n glob[c] = acc;\n }\n float centered0 = glob[0] - translations[row_i*3 + 0];\n float centered1 = glob[1] - translations[row_i*3 + 1];\n float centered2 = glob[2] - translations[row_i*3 + 2];\n float local = centered0 * rotations[(row_i*3 + ld)*3 + 0] + centered1 * rotations[(row_i*3 + ld)*3 + 1] + centered2 * rotations[(row_i*3 + ld)*3 + 2];\n cat[row_i * out_width + H*dhead + h * VP * 3 + tid] = local;\n }\n}\n\n__global__ void out_proj_kernel(\n const float* __restrict__ cat,\n const float* __restrict__ out_w,\n const float* __restrict__ out_b,\n float* __restrict__ out,\n int BN, int D, int INW)\n{\n int row = blockIdx.x;\n int od = threadIdx.x;\n if (od < D) {\n float acc = out_b[od];\n const float* w = out_w + od * INW;\n const float* x = cat + row * INW;\n for (int f = 0; f < INW; ++f) acc += x[f] * w[f];\n out[row * D + od] = acc;\n }\n}\n\ntorch::Tensor ipa_forward(\n torch::Tensor node_repr, torch::Tensor pair_repr,\n torch::Tensor translations, torch::Tensor rotations,\n torch::Tensor ln_w, torch::Tensor ln_b,\n torch::Tensor q_w, torch::Tensor k_w, torch::Tensor v_w,\n torch::Tensor qpts_w, torch::Tensor kpts_w, torch::Tensor vpts_w,\n torch::Tensor pair_w, torch::Tensor head_w,\n torch::Tensor out_w, torch::Tensor out_b,\n int64_t num_heads, int64_t num_query_points, int64_t num_value_points)\n{\n int B = (int)node_repr.size(0);\n int N = (int)node_repr.size(1);\n int D = (int)node_repr.size(2);\n int PD = (int)pair_repr.size(3);\n int H = (int)num_heads;\n int QP = (int)num_query_points;\n int VP = (int)num_value_points;\n int BN = B * N;\n int dhead = D / H;\n auto opts = node_repr.options();\n auto q = torch::empty({BN, H*dhead}, opts);\n auto k = torch::empty({BN, H*dhead}, opts);\n auto v = torch::empty({BN, H*dhead}, opts);\n auto qpl = torch::empty({BN, H*QP*3}, opts);\n auto kpl = torch::empty({BN, H*QP*3}, opts);\n auto vpl = torch::empty({BN, H*VP*3}, opts);\n auto cat = torch::empty({BN, H*dhead + H*VP*3}, opts);\n auto out = torch::empty({B, N, D}, opts);\n\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n ln_proj_kernel<<<BN, 128, 0, stream>>>(\n node_repr.data_ptr<float>(), ln_w.data_ptr<float>(), ln_b.data_ptr<float>(),\n q_w.data_ptr<float>(), k_w.data_ptr<float>(), v_w.data_ptr<float>(),\n qpts_w.data_ptr<float>(), kpts_w.data_ptr<float>(), vpts_w.data_ptr<float>(),\n q.data_ptr<float>(), k.data_ptr<float>(), v.data_ptr<float>(),\n qpl.data_ptr<float>(), kpl.data_ptr<float>(), vpl.data_ptr<float>(),\n BN, D, H, QP, VP);\n attn_kernel<<<B*N*H, 64, 0, stream>>>(\n pair_repr.data_ptr<float>(), translations.data_ptr<float>(), rotations.data_ptr<float>(),\n q.data_ptr<float>(), k.data_ptr<float>(), v.data_ptr<float>(),\n qpl.data_ptr<float>(), kpl.data_ptr<float>(), vpl.data_ptr<float>(),\n pair_w.data_ptr<float>(), head_w.data_ptr<float>(), cat.data_ptr<float>(),\n B, N, D, PD, H, QP, VP);\n int inw = H*dhead + H*VP*3;\n out_proj_kernel<<<BN, 64, 0, stream>>>(cat.data_ptr<float>(), out_w.data_ptr<float>(), out_b.data_ptr<float>(), out.data_ptr<float>(), BN, D, inw);\n return out;\n}\n'''\n\n_ipa_ext = load_inline(\n name='ipa_fused_ext_v1',\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=['ipa_forward'],\n verbose=False,\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, node_dim, pair_dim, num_heads, num_query_points, num_value_points):\n super().__init__()\n self.num_heads = num_heads\n self.head_dim = node_dim // num_heads\n self.num_query_points = num_query_points\n self.num_value_points = num_value_points\n assert node_dim % num_heads == 0\n d = self.head_dim\n self.q_proj = nn.Linear(node_dim, num_heads * d, bias=False)\n self.k_proj = nn.Linear(node_dim, num_heads * d, bias=False)\n self.v_proj = nn.Linear(node_dim, num_heads * d, bias=False)\n self.q_pts = nn.Linear(node_dim, num_heads * num_query_points * 3, bias=False)\n self.k_pts = nn.Linear(node_dim, num_heads * num_query_points * 3, bias=False)\n self.v_pts = nn.Linear(node_dim, num_heads * num_value_points * 3, bias=False)\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def forward(self, node_repr, pair_repr, translations, rotations):\n return _ipa_ext.ipa_forward(\n node_repr, pair_repr, translations, rotations,\n self.layer_norm.weight, self.layer_norm.bias,\n self.q_proj.weight, self.k_proj.weight, self.v_proj.weight,\n self.q_pts.weight, self.k_pts.weight, self.v_pts.weight,\n self.pair_bias.weight, self.head_weights,\n self.out_proj.weight, self.out_proj.bias,\n self.num_heads, self.num_query_points, self.num_value_points)\n" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.05 μs Runtime stats: mean=0.05μs median=?μs std=0.00μs Numerical precision: max_abs_err=5.07e-07 mean_abs_err=8.10e-08 max_rel_err=5.74e-03 Memory: 33.69 MB peak (ref: 34.48 MB, ratio: 0.9771x) Kernel launches: 8 (ref: 309, fusion ratio: 38.625) SOL score: 0.0114 Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (324 lines, 12,487 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>

torch::Tensor ipa_forward(
    torch::Tensor node_repr, torch::Tensor pair_repr,
    torch::Tensor translations, torch::Tensor rotations,
    torch::Tensor ln_w, torch::Tensor ln_b,
    torch::Tensor q_w, torch::Tensor k_w, torch::Tensor v_w,
    torch::Tensor qpts_w, torch::Tensor kpts_w, torch::Tensor vpts_w,
    torch::Tensor pair_w, torch::Tensor head_w,
    torch::Tensor out_w, torch::Tensor out_b,
    int64_t num_heads, int64_t num_query_points, int64_t num_value_points);
'''

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

static __device__ __forceinline__ float softplus_dev(float x) {
    if (x > 20.0f) return x;
    if (x < -20.0f) return expf(x);
    return log1pf(expf(x));
}

__global__ void ln_proj_kernel(
    const float* __restrict__ node,
    const float* __restrict__ ln_w,
    const float* __restrict__ ln_b,
    const float* __restrict__ q_w,
    const float* __restrict__ k_w,
    const float* __restrict__ v_w,
    const float* __restrict__ qpts_w,
    const float* __restrict__ kpts_w,
    const float* __restrict__ vpts_w,
    float* __restrict__ q,
    float* __restrict__ k,
    float* __restrict__ v,
    float* __restrict__ qpl,
    float* __restrict__ kpl,
    float* __restrict__ vpl,
    int BN, int D, int H, int QP, int VP)
{
    int row = blockIdx.x;
    int tid = threadIdx.x;
    __shared__ float red[128];
    __shared__ float xn[128];

    float s = 0.0f;
    for (int i = tid; i < D; i += blockDim.x) s += node[row * D + i];
    red[tid] = s;
    __syncthreads();
    for (int stride = blockDim.x/2; stride > 0; stride >>= 1) {
        if (tid < stride) red[tid] += red[tid + stride];
        __syncthreads();
    }
    float mean = red[0] / (float)D;

    s = 0.0f;
    for (int i = tid; i < D; i += blockDim.x) {
        float a = node[row * D + i] - mean;
        s += a * a;
    }
    red[tid] = s;
    __syncthreads();
    for (int stride = blockDim.x/2; stride > 0; stride >>= 1) {
        if (tid < stride) red[tid] += red[tid + stride];
        __syncthreads();
    }
    float inv_std = rsqrtf(red[0] / (float)D + 1.0e-5f);
    for (int i = tid; i < D; i += blockDim.x) {
        xn[i] = (node[row * D + i] - mean) * inv_std * ln_w[i] + ln_b[i];
    }
    __syncthreads();

    int dhead = D / H;
    int qkv = H * dhead;
    int qpo = H * QP * 3;
    int vpo = H * VP * 3;
    int total = 3*qkv + 2*qpo + vpo;

    for (int o = tid; o < total; o += blockDim.x) {
        const float* W;
        float* O;
        int oo, out_stride;
        if (o < qkv) { W = q_w; O = q; oo = o; out_stride = qkv; }
        else if (o < 2*qkv) { W = k_w; O = k; oo = o - qkv; out_stride = qkv; }
        else if (o < 3*qkv) { W = v_w; O = v; oo = o - 2*qkv; out_stride = qkv; }
        else if (o < 3*qkv + qpo) { W = qpts_w; O = qpl; oo = o - 3*qkv; out_stride = qpo; }
        else if (o < 3*qkv + 2*qpo) { W = kpts_w; O = kpl; oo = o - 3*qkv - qpo; out_stride = qpo; }
        else { W = vpts_w; O = vpl; oo = o - 3*qkv - 2*qpo; out_stride = vpo; }
        float acc = 0.0f;
        const float* wrow = W + oo * D;
        #pragma unroll
        for (int ii = 0; ii < 64; ++ii) {
            if (ii < D) acc += xn[ii] * wrow[ii];
        }
        O[row * out_stride + oo] = acc;
    }
}

__global__ void attn_kernel(
    const float* __restrict__ pair_repr,
    const float* __restrict__ translations,
    const float* __restrict__ rotations,
    const float* __restrict__ q,
    const float* __restrict__ k,
    const float* __restrict__ v,
    const float* __restrict__ qpl,
    const float* __restrict__ kpl,
    const float* __restrict__ vpl,
    const float* __restrict__ pair_w,
    const float* __restrict__ head_w,
    float* __restrict__ cat,
    int B, int N, int D, int PD, int H, int QP, int VP)
{
    int bid = blockIdx.x;
    int h = bid % H;
    int i = (bid / H) % N;
    int b = bid / (H * N);
    int tid = threadIdx.x;
    int dhead = D / H;
    int qkv_stride = H * dhead;
    int qp_stride = H * QP * 3;
    int vp_stride = H * VP * 3;
    int row_i = b * N + i;

    __shared__ float logits[64];
    __shared__ float attn[64];

    if (tid < N) {
        int j = tid;
        int row_j = b * N + j;
        float sc = 0.0f;
        for (int dd = 0; dd < dhead; ++dd) {
            sc += q[row_i * qkv_stride + h * dhead + dd] * k[row_j * qkv_stride + h * dhead + dd];
        }
        sc *= rsqrtf((float)dhead);

        float dist = 0.0f;
        for (int p = 0; p < QP; ++p) {
            float ql0 = qpl[row_i * qp_stride + (h*QP + p)*3 + 0];
            float ql1 = qpl[row_i * qp_stride + (h*QP + p)*3 + 1];
            float ql2 = qpl[row_i * qp_stride + (h*QP + p)*3 + 2];
            float kl0 = kpl[row_j * qp_stride + (h*QP + p)*3 + 0];
            float kl1 = kpl[row_j * qp_stride + (h*QP + p)*3 + 1];
            float kl2 = kpl[row_j * qp_stride + (h*QP + p)*3 + 2];
            for (int c = 0; c < 3; ++c) {
                float qg = ql0 * rotations[(row_i*3 + 0)*3 + c] + ql1 * rotations[(row_i*3 + 1)*3 + c] + ql2 * rotations[(row_i*3 + 2)*3 + c] + translations[row_i*3 + c];
                float kg = kl0 * rotations[(row_j*3 + 0)*3 + c] + kl1 * rotations[(row_j*3 + 1)*3 + c] + kl2 * rotations[(row_j*3 + 2)*3 + c] + translations[row_j*3 + c];
                float diff = qg - kg;
                dist += diff * diff;
            }
        }
        float pb = 0.0f;
        const float* pr = pair_repr + ((b * N + i) * N + j) * PD;
        const float* pw = pair_w + h * PD;
        for (int pp = 0; pp < PD; ++pp) pb += pr[pp] * pw[pp];
        float w = softplus_dev(head_w[h]);
        logits[j] = sc + pb - 0.5f * w * dist;
    }
    __syncthreads();

    if (tid == 0) {
        float m = logits[0];
        for (int j = 1; j < N; ++j) m = fmaxf(m, logits[j]);
        float s = 0.0f;
        for (int j = 0; j < N; ++j) {
            float e = expf(logits[j] - m);
            attn[j] = e;
            s += e;
        }
        float invs = 1.0f / s;
        for (int j = 0; j < N; ++j) attn[j] *= invs;
    }
    __syncthreads();

    int out_width = H * dhead + H * VP * 3;
    if (tid < dhead) {
        float acc = 0.0f;
        for (int j = 0; j < N; ++j) {
            int row_j = b * N + j;
            acc += attn[j] * v[row_j * qkv_stride + h * dhead + tid];
        }
        cat[row_i * out_width + h * dhead + tid] = acc;
    }
    if (tid < VP * 3) {
        int vp = tid / 3;
        int ld = tid - vp * 3;
        float glob[3];
        for (int c = 0; c < 3; ++c) {
            float acc = 0.0f;
            for (int j = 0; j < N; ++j) {
                int row_j = b * N + j;
                float vl0 = vpl[row_j * vp_stride + (h*VP + vp)*3 + 0];
                float vl1 = vpl[row_j * vp_stride + (h*VP + vp)*3 + 1];
                float vl2 = vpl[row_j * vp_stride + (h*VP + vp)*3 + 2];
                float vg = vl0 * rotations[(row_j*3 + 0)*3 + c] + vl1 * rotations[(row_j*3 + 1)*3 + c] + vl2 * rotations[(row_j*3 + 2)*3 + c] + translations[row_j*3 + c];
                acc += attn[j] * vg;
            }
            glob[c] = acc;
        }
        float centered0 = glob[0] - translations[row_i*3 + 0];
        float centered1 = glob[1] - translations[row_i*3 + 1];
        float centered2 = glob[2] - translations[row_i*3 + 2];
        float local = centered0 * rotations[(row_i*3 + ld)*3 + 0] + centered1 * rotations[(row_i*3 + ld)*3 + 1] + centered2 * rotations[(row_i*3 + ld)*3 + 2];
        cat[row_i * out_width + H*dhead + h * VP * 3 + tid] = local;
    }
}

__global__ void out_proj_kernel(
    const float* __restrict__ cat,
    const float* __restrict__ out_w,
    const float* __restrict__ out_b,
    float* __restrict__ out,
    int BN, int D, int INW)
{
    int row = blockIdx.x;
    int od = threadIdx.x;
    if (od < D) {
        float acc = out_b[od];
        const float* w = out_w + od * INW;
        const float* x = cat + row * INW;
        for (int f = 0; f < INW; ++f) acc += x[f] * w[f];
        out[row * D + od] = acc;
    }
}

torch::Tensor ipa_forward(
    torch::Tensor node_repr, torch::Tensor pair_repr,
    torch::Tensor translations, torch::Tensor rotations,
    torch::Tensor ln_w, torch::Tensor ln_b,
    torch::Tensor q_w, torch::Tensor k_w, torch::Tensor v_w,
    torch::Tensor qpts_w, torch::Tensor kpts_w, torch::Tensor vpts_w,
    torch::Tensor pair_w, torch::Tensor head_w,
    torch::Tensor out_w, torch::Tensor out_b,
    int64_t num_heads, int64_t num_query_points, int64_t num_value_points)
{
    int B = (int)node_repr.size(0);
    int N = (int)node_repr.size(1);
    int D = (int)node_repr.size(2);
    int PD = (int)pair_repr.size(3);
    int H = (int)num_heads;
    int QP = (int)num_query_points;
    int VP = (int)num_value_points;
    int BN = B * N;
    int dhead = D / H;
    auto opts = node_repr.options();
    auto q = torch::empty({BN, H*dhead}, opts);
    auto k = torch::empty({BN, H*dhead}, opts);
    auto v = torch::empty({BN, H*dhead}, opts);
    auto qpl = torch::empty({BN, H*QP*3}, opts);
    auto kpl = torch::empty({BN, H*QP*3}, opts);
    auto vpl = torch::empty({BN, H*VP*3}, opts);
    auto cat = torch::empty({BN, H*dhead + H*VP*3}, opts);
    auto out = torch::empty({B, N, D}, opts);

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    ln_proj_kernel<<<BN, 128, 0, stream>>>(
        node_repr.data_ptr<float>(), ln_w.data_ptr<float>(), ln_b.data_ptr<float>(),
        q_w.data_ptr<float>(), k_w.data_ptr<float>(), v_w.data_ptr<float>(),
        qpts_w.data_ptr<float>(), kpts_w.data_ptr<float>(), vpts_w.data_ptr<float>(),
        q.data_ptr<float>(), k.data_ptr<float>(), v.data_ptr<float>(),
        qpl.data_ptr<float>(), kpl.data_ptr<float>(), vpl.data_ptr<float>(),
        BN, D, H, QP, VP);
    attn_kernel<<<B*N*H, 64, 0, stream>>>(
        pair_repr.data_ptr<float>(), translations.data_ptr<float>(), rotations.data_ptr<float>(),
        q.data_ptr<float>(), k.data_ptr<float>(), v.data_ptr<float>(),
        qpl.data_ptr<float>(), kpl.data_ptr<float>(), vpl.data_ptr<float>(),
        pair_w.data_ptr<float>(), head_w.data_ptr<float>(), cat.data_ptr<float>(),
        B, N, D, PD, H, QP, VP);
    int inw = H*dhead + H*VP*3;
    out_proj_kernel<<<BN, 64, 0, stream>>>(cat.data_ptr<float>(), out_w.data_ptr<float>(), out_b.data_ptr<float>(), out.data_ptr<float>(), BN, D, inw);
    return out;
}
'''

_ipa_ext = load_inline(
    name='ipa_fused_ext_v1',
    cpp_sources=_cpp_src,
    cuda_sources=_cuda_src,
    functions=['ipa_forward'],
    verbose=False,
    extra_cflags=['-O3'],
    extra_cuda_cflags=['-O3']
)

class ModelNew(nn.Module):
    def __init__(self, node_dim, pair_dim, num_heads, num_query_points, num_value_points):
        super().__init__()
        self.num_heads = num_heads
        self.head_dim = node_dim // num_heads
        self.num_query_points = num_query_points
        self.num_value_points = num_value_points
        assert node_dim % num_heads == 0
        d = self.head_dim
        self.q_proj = nn.Linear(node_dim, num_heads * d, bias=False)
        self.k_proj = nn.Linear(node_dim, num_heads * d, bias=False)
        self.v_proj = nn.Linear(node_dim, num_heads * d, bias=False)
        self.q_pts = nn.Linear(node_dim, num_heads * num_query_points * 3, bias=False)
        self.k_pts = nn.Linear(node_dim, num_heads * num_query_points * 3, bias=False)
        self.v_pts = nn.Linear(node_dim, num_heads * num_value_points * 3, bias=False)
        self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)
        self.head_weights = nn.Parameter(torch.zeros(num_heads))
        out_features = num_heads * d + num_heads * num_value_points * 3
        self.out_proj = nn.Linear(out_features, node_dim)
        self.layer_norm = nn.LayerNorm(node_dim)

    def forward(self, node_repr, pair_repr, translations, rotations):
        return _ipa_ext.ipa_forward(
            node_repr, pair_repr, translations, rotations,
            self.layer_norm.weight, self.layer_norm.bias,
            self.q_proj.weight, self.k_proj.weight, self.v_proj.weight,
            self.q_pts.weight, self.k_pts.weight, self.v_pts.weight,
            self.pair_bias.weight, self.head_weights,
            self.out_proj.weight, self.out_proj.bias,
            self.num_heads, self.num_query_points, self.num_value_points)