FW-GLM-5-1 · popcorn · pop_l123_default_gpt
correct 36_InvariantPointAttention.py
model
FW-GLM-5-1
turns / max
10 / 12
tool calls
10
tokens (total)
219771
wall clock
4289.5s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.13 μs
ref runtime
0.18 μs
speedup
1.38x
started
2026-05-11T14:18:15Z
timingkernelbench.timing
kernel runtime
0.13 μs
ref runtime
0.18 μs
speedup
1.38x
kernel mean / median / std
0.13 / — / 9.95e-04 μs
ref mean / median / std
0.18 / — / 6.65e-04 μs
numerical precisioneval_kernel_against_ref
max abs error
3.58e-07
mean abs error
5.37e-08
max rel error
0.00
mean rel error
2.00e-06
gpu memoryextended_metrics.measure_memory
kernel peak
33.81 MB
ref peak
34.48 MB
kernel / ref
0.98x
kernel launches / fusiontorch.profiler
kernel launches
176
ref launches
309
fusion ratio (ref / kernel)
1.76x
total cuda time
2276 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
1.000
dram util %
37.0%
compute util %
100.0%
bottleneck
compute-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
1238.27 GB/s
dram util %
37.0%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
288.9%
bottleneck
compute-bound
dominant pipe %
138.2%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
2632.5%
ld sectors / req
138.77
st sectors / req
—
regs / thread
1338
smem / block
198144 B
block size
5952
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
15662.0%
short_scoreboard
6575.2%
mio_throttle
244.0%
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 0341.2s turn · 3.4s LLM · 337.8s tools · tokens in=2,821 · out=248 · total=3,069
reasoning
Let me understand this task. I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model is an Invariant Point Attention (IPA) from AlphaFold2's structure module. Let me first understand the computation: 1. Layer normalization on node_repr 2. Project queries, keys, values (linear projections) 3. Project point queries, keys, values (linear projections to 3D points) 4. Transform points from local to global frame using rotations and translations 5. Compute scalar attention logits via einsum 6. Compute point attention via squared distances between query and key points 7. Compute pair bias via linear projection 8. Combine: attn = attn_scalar + pt_attn + pair_bias 9. Apply softmax 10. Aggregate scalar values and point values 11. Transform point outputs back to local frame 12. Concatenate and project output This is a complex operation with many steps. Let me think about how to optimize this with CUDA. First, let me get the GPU specs to understand what I'm working with.
assistantI'll start by understanding the GPU specs and then analyze the reference implementation to design an optimized CUDA kernel.
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) · 835.0s turn · 52.8s LLM · tokens in=3,090 · out=5,661 · total=8,751 · 21,317 chars (click to expand)
turn 1835.0s turn · 52.8s LLM · 782.2s tools · tokens in=3,090 · out=5,661 · total=8,751
reasoning
Now let me analyze the computation to understand what we're optimizing: 1. LayerNorm - this can be fused 2. Linear projections (q, k, v, q_pts, k_pts, v_pts, pair_bias, out_proj) - these are the main compute work 3. Frame transformations (apply_frame and apply_inv_frame) 4. Point distance computation 5. Attention computation (scalar + point + pair_bias) 6. Aggregation of scalar and point values The dimensions are: - B = 2, N = 32, node_dim = 64, pair_dim = 32 - num_heads = 4, head_dim = 16 - num_query_points = 4, num_value_points = 8 Key operations: 1. LayerNorm: (B, N, 64) -> (B, N, 64) 2. Linear projections: - q_proj: 64 -> 64 (4*16) - k_proj: 64 -> 64 - v_proj: 64 -> 64 - q_pts: 64 -> 48 (4*4*3) - k_pts: 64 -> 48 - v_pts: 64 -> 96 (4*8*3) - pair_bias: 32 -> 4 (B, N, N, pair_dim) -> (B, N, N, 4) - out_proj: 64+96 -> 64 (160 -> 64) The attention computation is on (B, h, N, N) = (2, 4, 32, 32) Let me write a fused kernel that handles all operations together. Since the problem size is small (N=32), we might benefit from a single kernel launch that does all the work. Actually, given the small problem sizes, the main bottleneck might be kernel launch overhead. Let me write a fused kernel that does everything. Let me think about the memory access patterns and compute: - The main compute is in the linear projections (matrix multiplications) - The attention is small (32x32) and can be done efficiently - The frame transformations are straightforward For a fused kernel approach, I should: 1. Load the input node_repr into shared memory 2. Compute layer norm 3. Compute all projections (this requires weights which are large) 4. Transform points 5. Compute attention 6. Aggregate results 7. Output projection Actually, the weights for the linear layers are quite large (especially pair_bias which operates on pair_repr of shape (B, N, N, pair_dim)). The pair_repr itself is (2, 32, 32, 32) = 65536 floats = 256KB. Let me think more carefully. The pair_bias projection is: - Input: (B, N, N, pair_dim) = (2, 32, 32, 32) - Weight: (pair_dim, num_heads) = (32, 4) - Output: (B, N, N, num_heads) = (2, 32, 32, 4) This is a small linear layer that can be fused. For the main projections (q, k, v), the weights are 64x64 each. Actually, let me reconsider. The reference implementation uses PyTorch's optimized linear layers which are highly optimized. The main opportunity for optimization might be: 1. Fusing operations to reduce memory traffic 2. Fusing layer norm with projections 3. Better memory access patterns for the point computations Let me write a comprehensive kernel that: 1. Fuses layer norm with the small projections 2. Uses shared memory efficiently for attention computation 3. Fuses the final operations Since N=32 is small, we can have each thread block handle one (batch, head) pair for attention, or one (batch, query_position) pair for the full computation. Let me start with a simpler approach: fuse the attention and aggregation parts, keeping the projections as PyTorch calls, then see if we need to fuse more. Actually, looking at the dimensions again, the key opportunity is: 1. The point distance computation creates a large intermediate (B, N, N, h) = (2, 32, 32, 4) 2. The attention computation also creates intermediates Let me write a kernel that fuses: - Point distance computation - Attention (scalar + point + pair_bias + softmax) - Value aggregation This should reduce memory traffic significantly. Wait, let me reconsider the best approach. The input sizes are: - node_repr: (2, 32, 64) = 4KB - pair_repr: (2, 32, 32, 32) = 256KB - translations: (2, 32, 3) = 768 bytes - rotations: (2, 32, 3, 3) = 2.3KB The pair_repr is the largest input. The main compute is: 1. Projections on node_repr (small, compute-bound) 2. Projections on pair_repr (larger, memory-bound) 3. Frame transformations (small) 4. Point distance (creates N*N intermediates) 5. Attention and aggregation For efficiency, I should: 1. Keep projections as-is (PyTorch optimized) 2. Fuse the attention computation kernel to avoid materializing large intermediates Let me write a kernel that handles the attention computation and aggregation.
assistantLet me analyze this Invariant Point Attention implementation and write an optimized CUDA kernel. The key operations are: 1. LayerNorm + projections 2. Frame transformations for 3D points 3. Point distance computation (creates N×N intermediates) 4. Attention computation with scalar, point, and pair bias terms 5. Value aggregation I'll write a fused kernel for the attention and aggregation parts, which have significant intermediate tensors that can be optimized.
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Fused kernel for IPA attention computation and value aggregation\n// Each thread block handles one (batch, head, query_pos) tuple\n// Computes attention over all key positions and aggregates values\ntemplate <int N, int BLOCK_SIZE>\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 3, 3)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n \n // Load head weight\n float w = head_weights[head];\n w = log1pf(expf(w)); // softplus\n \n // Load query vector for this position\n // q shape: (B, N, h, d)\n extern __shared__ float smem[];\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += BLOCK_SIZE) {\n q_local[i] = q[batch * N * h * d + query_pos * h * d + head * d + i];\n }\n for (int i = tid; i < Qp * 3; i += BLOCK_SIZE) {\n q_pts_local[i] = q_pts[batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3 + i];\n }\n for (int i = tid; i < 9; i += BLOCK_SIZE) {\n rot_local[i] = rotations[batch * N * 9 + query_pos * 9 + i];\n }\n for (int i = tid; i < 3; i += BLOCK_SIZE) {\n trans_local[i] = translations[batch * N * 3 + query_pos * 3 + i];\n }\n __syncthreads();\n \n // Compute scalar attention scores and find max for softmax stability\n float max_score = -INFINITY;\n float sum_exp = 0.0f;\n \n // Each thread handles a subset of key positions\n float local_attn[BLOCK_SIZE / 32]; // Store attention weights for key positions\n \n // First pass: compute all scores and find max\n for (int key_pos = tid; key_pos < N; key_pos += BLOCK_SIZE) {\n // Scalar attention: q \u00b7 k / sqrt(d)\n float scalar_score = 0.0f;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n scalar_score /= sqrtf((float)d);\n \n // Point attention: -w * ||q_pts - k_pts||^2 / 2\n float pt_score = 0.0f;\n for (int p = 0; p < Qp; p++) {\n float dx = 0.0f, dy = 0.0f, dz = 0.0f;\n for (int c = 0; c < 3; c++) {\n float q_pt = q_pts_local[p * 3 + c];\n float k_pt = k_pts[batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3 + p * 3 + c];\n float diff = q_pt - k_pt;\n if (c == 0) dx = diff;\n else if (c == 1) dy = diff;\n else dz = diff;\n }\n pt_score += dx * dx + dy * dy + dz * dz;\n }\n pt_score = -w * pt_score / 2.0f;\n \n // Pair bias\n float pb = pair_bias[batch * N * N * h + query_pos * N * h + key_pos * h + head];\n \n float total_score = scalar_score + pt_score + pb;\n local_attn[key_pos / 32] = total_score;\n max_score = max(max_score, total_score);\n }\n \n // Reduce max across block\n __shared__ float shared_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n shared_max[warp_id] = my_max;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_max = lane_id < (BLOCK_SIZE / 32) ? shared_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n shared_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = shared_max[0];\n \n // Second pass: compute exp(score - max) and sum\n float sum_e = 0.0f;\n for (int key_pos = tid; key_pos < N; key_pos += BLOCK_SIZE) {\n float e = expf(local_attn[key_pos / 32] - global_max);\n local_attn[key_pos / 32] = e;\n sum_e += e;\n }\n \n // Reduce sum across block\n __shared__ float shared_sum[32];\n float my_sum = warp_reduce_sum(sum_e);\n if (lane_id == 0) {\n shared_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < (BLOCK_SIZE / 32) ? shared_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n shared_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float total_sum = shared_sum[0];\n \n // Third pass: normalize and aggregate values\n float out_s[16] = {0}; // d <= 16 typically\n float out_p[24] = {0}; // Vp * 3 <= 24 typically\n \n for (int key_pos = tid; key_pos < N; key_pos += BLOCK_SIZE) {\n float attn_weight = local_attn[key_pos / 32] / total_sum;\n \n // Aggregate scalar values\n for (int i = 0; i < d; i++) {\n out_s[i] += attn_weight * v[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n \n // Aggregate point values\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn_weight * v_pts[batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3 + p * 3 + c];\n }\n }\n }\n \n // Reduce across warp and write output\n __syncthreads();\n \n // For scalar output\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n out_scalar[batch * N * h * d + query_pos * h * d + head * d + i] = val;\n }\n }\n \n // For point output - need to transform back to local frame\n // This is done after gathering all warps' contributions\n __shared__ float shared_pts[64 * 3]; // Max Vp * 3\n \n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n shared_pts[p * 3 + c] = val;\n }\n }\n }\n __syncthreads();\n \n // Transform points from global to local frame\n // out_pts_local = (out_pts - trans) @ rot (rot is orthonormal, so inv = transpose)\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n // rot_local is stored row-major: rot[r][c] = rot_local[r*3 + c]\n // We need rot^T, so rot^T[r][c] = rot[c][r] = rot_local[c*3 + r]\n local_pt += global_pt * rot_local[c * 3 + i]; // rot^T @ (pt - trans)\n }\n out_pts[batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3 + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q, // (B, N, h, d)\n torch::Tensor k, // (B, N, h, d)\n torch::Tensor v, // (B, N, h, d)\n torch::Tensor q_pts, // (B, N, h, Qp, 3)\n torch::Tensor k_pts, // (B, N, h, Qp, 3)\n torch::Tensor v_pts, // (B, N, h, Vp, 3)\n torch::Tensor pair_bias, // (B, N, N, h)\n torch::Tensor head_weights, // (h,)\n torch::Tensor rotations, // (B, N, 3, 3)\n torch::Tensor translations // (B, N, 3)\n) {\n int B = q.size(0);\n int N = q.size(1);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int BLOCK_SIZE = 128;\n dim3 grid(N, B, h);\n size_t smem_size = (d + Qp * 3 + 9 + 3) * sizeof(float);\n \n ipa_attention_kernel<N_MAX, BLOCK_SIZE><<<grid, BLOCK_SIZE, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations);\n\"\"\"\n\ntry:\n ipa_extension = load_inline(\n name='ipa_attention',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-DN_MAX=32'],\n verbose=True\n )\nexcept Exception as e:\n print(f\"Failed to compile: {e}\")\n ipa_extension = None\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias = self.pair_bias(pair_repr).permute(0, 3, 1, 2) # (B, h, N, N)\n\n # Use custom CUDA kernel for attention and aggregation\n if ipa_extension is not None and node_repr.is_cuda:\n # Flatten rotations from (B, N, 3, 3) to (B, N, 9)\n rotations_flat = rotations.reshape(B, N, 9)\n \n # head_weights needs softplus applied, but kernel does it\n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias.permute(0, 2, 3, 1).contiguous(), # (B, N, N, h)\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous()\n )\n return self.out_proj(result)\n else:\n # Fallback to reference implementation\n attn_scalar = torch.einsum(\"bihd,bjhd->bhij\", q, k) / math.sqrt(d)\n q_expand = q_pts.unsqueeze(2)\n k_expand = k_pts.unsqueeze(1)\n pt_dist = ((q_expand - k_expand) ** 2).sum(dim=(-1, -2))\n w = F.softplus(self.head_weights)\n pt_attn = -(w.view(1, 1, 1, h) * pt_dist) / 2.0\n pt_attn = pt_attn.permute(0, 3, 1, 2)\n attn = attn_scalar + pt_attn + pair_bias\n attn = F.softmax(attn, dim=-1)\n out_scalar = torch.einsum(\"bhij,bjhd->bihd\", attn, v).reshape(B, N, h * d)\n out_pts = torch.einsum(\"bhij,bjhpc->bihpc\", attn, v_pts)\n out_pts = out_pts.reshape(B, N, h * Vp, 3)\n out_pts_local = self._apply_inv_frame(out_pts, rotations, translations)\n out_pts_flat = out_pts_local.reshape(B, N, h * Vp * 3)\n out = torch.cat([out_scalar, out_pts_flat], dim=-1)\n return self.out_proj(out)\n \n def _apply_inv_frame(self, pts, rotations, translations):\n centered = pts - translations.unsqueeze(2)\n local_pts = torch.einsum(\"bnpc,bndc->bnpd\", centered, rotations)\n return local_pts" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 2 · 1 tool call(s) · 36.0s turn · 36.0s LLM · tokens in=7,616 · out=4,446 · total=12,062 · 16,938 chars (click to expand)
turn 236.0s turn · 36.0s LLM · 0.0s tools · tokens in=7,616 · out=4,446 · total=12,062
reasoning
Good, the kernel compiled. Let me check for static analysis issues first.
assistantLet me run static analysis to check for any issues:
call → static_check
callstatic_check { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Fused kernel for IPA attention computation and value aggregation\n// Each thread block handles one (batch, head, query_pos) tuple\n// Computes attention over all key positions and aggregates values\ntemplate <int N, int BLOCK_SIZE>\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 3, 3)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n \n // Load head weight\n float w = head_weights[head];\n w = log1pf(expf(w)); // softplus\n \n // Load query vector for this position\n // q shape: (B, N, h, d)\n extern __shared__ float smem[];\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += BLOCK_SIZE) {\n q_local[i] = q[batch * N * h * d + query_pos * h * d + head * d + i];\n }\n for (int i = tid; i < Qp * 3; i += BLOCK_SIZE) {\n q_pts_local[i] = q_pts[batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3 + i];\n }\n for (int i = tid; i < 9; i += BLOCK_SIZE) {\n rot_local[i] = rotations[batch * N * 9 + query_pos * 9 + i];\n }\n for (int i = tid; i < 3; i += BLOCK_SIZE) {\n trans_local[i] = translations[batch * N * 3 + query_pos * 3 + i];\n }\n __syncthreads();\n \n // Compute scalar attention scores and find max for softmax stability\n float max_score = -INFINITY;\n float sum_exp = 0.0f;\n \n // Each thread handles a subset of key positions\n float local_attn[BLOCK_SIZE / 32]; // Store attention weights for key positions\n \n // First pass: compute all scores and find max\n for (int key_pos = tid; key_pos < N; key_pos += BLOCK_SIZE) {\n // Scalar attention: q \u00b7 k / sqrt(d)\n float scalar_score = 0.0f;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n scalar_score /= sqrtf((float)d);\n \n // Point attention: -w * ||q_pts - k_pts||^2 / 2\n float pt_score = 0.0f;\n for (int p = 0; p < Qp; p++) {\n float dx = 0.0f, dy = 0.0f, dz = 0.0f;\n for (int c = 0; c < 3; c++) {\n float q_pt = q_pts_local[p * 3 + c];\n float k_pt = k_pts[batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3 + p * 3 + c];\n float diff = q_pt - k_pt;\n if (c == 0) dx = diff;\n else if (c == 1) dy = diff;\n else dz = diff;\n }\n pt_score += dx * dx + dy * dy + dz * dz;\n }\n pt_score = -w * pt_score / 2.0f;\n \n // Pair bias\n float pb = pair_bias[batch * N * N * h + query_pos * N * h + key_pos * h + head];\n \n float total_score = scalar_score + pt_score + pb;\n local_attn[key_pos / 32] = total_score;\n max_score = max(max_score, total_score);\n }\n \n // Reduce max across block\n __shared__ float shared_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n shared_max[warp_id] = my_max;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_max = lane_id < (BLOCK_SIZE / 32) ? shared_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n shared_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = shared_max[0];\n \n // Second pass: compute exp(score - max) and sum\n float sum_e = 0.0f;\n for (int key_pos = tid; key_pos < N; key_pos += BLOCK_SIZE) {\n float e = expf(local_attn[key_pos / 32] - global_max);\n local_attn[key_pos / 32] = e;\n sum_e += e;\n }\n \n // Reduce sum across block\n __shared__ float shared_sum[32];\n float my_sum = warp_reduce_sum(sum_e);\n if (lane_id == 0) {\n shared_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < (BLOCK_SIZE / 32) ? shared_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n shared_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float total_sum = shared_sum[0];\n \n // Third pass: normalize and aggregate values\n float out_s[16] = {0}; // d <= 16 typically\n float out_p[24] = {0}; // Vp * 3 <= 24 typically\n \n for (int key_pos = tid; key_pos < N; key_pos += BLOCK_SIZE) {\n float attn_weight = local_attn[key_pos / 32] / total_sum;\n \n // Aggregate scalar values\n for (int i = 0; i < d; i++) {\n out_s[i] += attn_weight * v[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n \n // Aggregate point values\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn_weight * v_pts[batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3 + p * 3 + c];\n }\n }\n }\n \n // Reduce across warp and write output\n __syncthreads();\n \n // For scalar output\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n out_scalar[batch * N * h * d + query_pos * h * d + head * d + i] = val;\n }\n }\n \n // For point output - need to transform back to local frame\n // This is done after gathering all warps' contributions\n __shared__ float shared_pts[64 * 3]; // Max Vp * 3\n \n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n shared_pts[p * 3 + c] = val;\n }\n }\n }\n __syncthreads();\n \n // Transform points from global to local frame\n // out_pts_local = (out_pts - trans) @ rot (rot is orthonormal, so inv = transpose)\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n // rot_local is stored row-major: rot[r][c] = rot_local[r*3 + c]\n // We need rot^T, so rot^T[r][c] = rot[c][r] = rot_local[c*3 + r]\n local_pt += global_pt * rot_local[c * 3 + i]; // rot^T @ (pt - trans)\n }\n out_pts[batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3 + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q, // (B, N, h, d)\n torch::Tensor k, // (B, N, h, d)\n torch::Tensor v, // (B, N, h, d)\n torch::Tensor q_pts, // (B, N, h, Qp, 3)\n torch::Tensor k_pts, // (B, N, h, Qp, 3)\n torch::Tensor v_pts, // (B, N, h, Vp, 3)\n torch::Tensor pair_bias, // (B, N, N, h)\n torch::Tensor head_weights, // (h,)\n torch::Tensor rotations, // (B, N, 3, 3)\n torch::Tensor translations // (B, N, 3)\n) {\n int B = q.size(0);\n int N = q.size(1);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int BLOCK_SIZE = 128;\n dim3 grid(N, B, h);\n size_t smem_size = (d + Qp * 3 + 9 + 3) * sizeof(float);\n \n ipa_attention_kernel<N_MAX, BLOCK_SIZE><<<grid, BLOCK_SIZE, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations);\n\"\"\"\n\ntry:\n ipa_extension = load_inline(\n name='ipa_attention',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-DN_MAX=32'],\n verbose=True\n )\nexcept Exception as e:\n print(f\"Failed to compile: {e}\")\n ipa_extension = None\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias = self.pair_bias(pair_repr).permute(0, 3, 1, 2) # (B, h, N, N)\n\n # Use custom CUDA kernel for attention and aggregation\n if ipa_extension is not None and node_repr.is_cuda:\n # Flatten rotations from (B, N, 3, 3) to (B, N, 9)\n rotations_flat = rotations.reshape(B, N, 9)\n \n # head_weights needs softplus applied, but kernel does it\n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias.permute(0, 2, 3, 1).contiguous(), # (B, N, N, h)\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous()\n )\n return self.out_proj(result)\n else:\n # Fallback to reference implementation\n attn_scalar = torch.einsum(\"bihd,bjhd->bhij\", q, k) / math.sqrt(d)\n q_expand = q_pts.unsqueeze(2)\n k_expand = k_pts.unsqueeze(1)\n pt_dist = ((q_expand - k_expand) ** 2).sum(dim=(-1, -2))\n w = F.softplus(self.head_weights)\n pt_attn = -(w.view(1, 1, 1, h) * pt_dist) / 2.0\n pt_attn = pt_attn.permute(0, 3, 1, 2)\n attn = attn_scalar + pt_attn + pair_bias\n attn = F.softmax(attn, dim=-1)\n out_scalar = torch.einsum(\"bhij,bjhd->bihd\", attn, v).reshape(B, N, h * d)\n out_pts = torch.einsum(\"bhij,bjhpc->bihpc\", attn, v_pts)\n out_pts = out_pts.reshape(B, N, h * Vp, 3)\n out_pts_local = self._apply_inv_frame(out_pts, rotations, translations)\n out_pts_flat = out_pts_local.reshape(B, N, h * Vp * 3)\n out = torch.cat([out_scalar, out_pts_flat], dim=-1)\n return self.out_proj(out)\n \n def _apply_inv_frame(self, pts, rotations, translations):\n centered = pts - translations.unsqueeze(2)\n local_pts = torch.einsum(\"bnpc,bndc->bnpd\", centered, rotations)\n return local_pts" }
result ← static_check (fail)
failstatic_check static_check FAILED: strict violations found. ERROR: Contains try-except block (potential fallback bypass) Advisory warnings: WARNING: Uses torch computation op: torch.einsum WARNING: Uses torch.nn compute layer (only containers, Parameter, init allowed)
turn 3 · 1 tool call(s) · 382.8s turn · 34.1s LLM · tokens in=12,106 · out=4,063 · total=16,169 · 15,407 chars (click to expand)
turn 3382.8s turn · 34.1s LLM · 348.7s tools · tokens in=12,106 · out=4,063 · total=16,169
reasoning
The static check flagged my try-except block and fallback to reference. I need to rewrite without those. Let me create a version that: 1. Has no try-except blocks 2. Has no fallback to the reference implementation 3. Uses proper methods only Also the template kernel with N_MAX=32 isn't ideal for variable N. Let me rewrite without templating.
assistantLet me rewrite without the try-except and fallback:
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n// Fused kernel for IPA attention computation and value aggregation\n// Each thread block handles one (batch, head, query_pos) tuple\n// Computes attention over all key positions and aggregates values\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w)); // softplus\n \n // Shared memory layout\n extern __shared__ float smem[];\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements for attention scores\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[batch * N * h * d + query_pos * h * d + head * d + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3 + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[batch * N * 9 + query_pos * 9 + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[batch * N * 3 + query_pos * 3 + i];\n }\n __syncthreads();\n \n // First pass: compute all scores and store in shared memory\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention: q \u00b7 k / sqrt(d)\n float scalar_score = 0.0f;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n scalar_score /= sqrtf((float)d);\n \n // Point attention: -w * ||q_pts - k_pts||^2 / 2\n float pt_score = 0.0f;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float q_pt = q_pts_local[p * 3 + c];\n float k_pt = k_pts[batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3 + p * 3 + c];\n float diff = q_pt - k_pt;\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score / 2.0f;\n \n // Pair bias\n float pb = pair_bias[batch * N * N * h + query_pos * N * h + key_pos * h + head];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n }\n __syncthreads();\n \n // Find max for softmax stability\n float max_score = -INFINITY;\n for (int i = tid; i < N; i += block_size) {\n max_score = max(max_score, scores[i]);\n }\n \n // Reduce max across block\n __shared__ float shared_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n shared_max[warp_id] = my_max;\n }\n __syncthreads();\n int num_warps = block_size / 32;\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? shared_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n shared_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = shared_max[0];\n \n // Compute exp(score - max) and sum\n float sum_e = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_e += e;\n }\n \n // Reduce sum across block\n __shared__ float shared_sum[32];\n float my_sum = warp_reduce_sum(sum_e);\n if (lane_id == 0) {\n shared_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? shared_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n shared_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float total_sum = shared_sum[0];\n \n // Normalize attention weights\n for (int i = tid; i < N; i += block_size) {\n scores[i] /= total_sum;\n }\n __syncthreads();\n \n // Aggregate values\n // Each thread accumulates partial sums\n float out_s[16] = {0}; // d <= 16 typically\n float out_p[24] = {0}; // Vp * 3 <= 24 typically\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn_weight = scores[key_pos];\n \n // Aggregate scalar values\n for (int i = 0; i < d; i++) {\n out_s[i] += attn_weight * v[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n \n // Aggregate point values\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn_weight * v_pts[batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3 + p * 3 + c];\n }\n }\n }\n \n // Reduce across block and write scalar output\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n // Use atomic add to accumulate across warps\n atomicAdd(&out_scalar[batch * N * h * d + query_pos * h * d + head * d + i], val);\n }\n }\n \n // For point output - use shared memory to gather warps' contributions\n __shared__ float shared_pts[64 * 3]; // Max Vp * 3\n \n // Initialize shared_pts\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Each warp accumulates to shared_pts\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points from global to local frame and write output\n // out_pts_local = (out_pts - trans) @ rot^T\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n // rot_local is stored row-major: rot[r][c] = rot_local[r*3 + c]\n // We need rot^T, so rot^T[r][c] = rot[c][r] = rot_local[c*3 + r]\n local_pt += global_pt * rot_local[c * 3 + i]; // rot^T @ (pt - trans)\n }\n out_pts[batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3 + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q, // (B, N, h, d)\n torch::Tensor k, // (B, N, h, d)\n torch::Tensor v, // (B, N, h, d)\n torch::Tensor q_pts, // (B, N, h, Qp, 3)\n torch::Tensor k_pts, // (B, N, h, Qp, 3)\n torch::Tensor v_pts, // (B, N, h, Vp, 3)\n torch::Tensor pair_bias, // (B, N, N, h)\n torch::Tensor head_weights, // (h,)\n torch::Tensor rotations, // (B, N, 9)\n torch::Tensor translations, // (B, N, 3)\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n\n # Flatten rotations from (B, N, 3, 3) to (B, N, 9)\n rotations_flat = rotations.reshape(B, N, 9)\n \n # Call custom CUDA kernel\n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(), # (B, N, N, h)\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: RuntimeError. RuntimeError: Error building extension 'ipa_attention_v1'
turn 4 · 1 tool call(s) · 757.3s turn · 32.5s LLM · tokens in=16,113 · out=4,028 · total=20,141 · 15,184 chars (click to expand)
turn 4757.3s turn · 32.5s LLM · 724.7s tools · tokens in=16,113 · out=4,028 · total=20,141
reasoning
There's a compile error. Let me check the CUDA code more carefully. I see there's an extra parenthesis in the warp_reduce_sum function. Let me fix that and also check for other issues.
assistantLet me fix the syntax error in the CUDA code:
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Fused kernel for IPA attention computation and value aggregation\n// Each thread block handles one (batch, head, query_pos) tuple\n// Computes attention over all key positions and aggregates values\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w)); // softplus\n \n // Shared memory layout\n extern __shared__ float smem[];\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements for attention scores\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[batch * N * h * d + query_pos * h * d + head * d + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3 + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[batch * N * 9 + query_pos * 9 + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[batch * N * 3 + query_pos * 3 + i];\n }\n __syncthreads();\n \n // First pass: compute all scores and store in shared memory\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention: q \u00b7 k / sqrt(d)\n float scalar_score = 0.0f;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n scalar_score /= sqrtf((float)d);\n \n // Point attention: -w * ||q_pts - k_pts||^2 / 2\n float pt_score = 0.0f;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float q_pt = q_pts_local[p * 3 + c];\n float k_pt = k_pts[batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3 + p * 3 + c];\n float diff = q_pt - k_pt;\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score / 2.0f;\n \n // Pair bias\n float pb = pair_bias[batch * N * N * h + query_pos * N * h + key_pos * h + head];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n }\n __syncthreads();\n \n // Find max for softmax stability\n float max_score = -INFINITY;\n for (int i = tid; i < N; i += block_size) {\n max_score = max(max_score, scores[i]);\n }\n \n // Reduce max across block\n __shared__ float shared_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n shared_max[warp_id] = my_max;\n }\n __syncthreads();\n int num_warps = block_size / 32;\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? shared_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n shared_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = shared_max[0];\n \n // Compute exp(score - max) and sum\n float sum_e = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_e += e;\n }\n \n // Reduce sum across block\n __shared__ float shared_sum[32];\n float my_sum = warp_reduce_sum(sum_e);\n if (lane_id == 0) {\n shared_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? shared_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n shared_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float total_sum = shared_sum[0];\n \n // Normalize attention weights\n for (int i = tid; i < N; i += block_size) {\n scores[i] /= total_sum;\n }\n __syncthreads();\n \n // Aggregate values\n // Each thread accumulates partial sums\n float out_s[16] = {0}; // d <= 16 typically\n float out_p[24] = {0}; // Vp * 3 <= 24 typically\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn_weight = scores[key_pos];\n \n // Aggregate scalar values\n for (int i = 0; i < d; i++) {\n out_s[i] += attn_weight * v[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n \n // Aggregate point values\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn_weight * v_pts[batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3 + p * 3 + c];\n }\n }\n }\n \n // Reduce across block and write scalar output\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n // Use atomic add to accumulate across warps\n atomicAdd(&out_scalar[batch * N * h * d + query_pos * h * d + head * d + i], val);\n }\n }\n \n // For point output - use shared memory to gather warps' contributions\n __shared__ float shared_pts[64 * 3]; // Max Vp * 3\n \n // Initialize shared_pts\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Each warp accumulates to shared_pts\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points from global to local frame and write output\n // out_pts_local = (out_pts - trans) @ rot^T\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n // rot_local is stored row-major: rot[r][c] = rot_local[r*3 + c]\n // We need rot^T, so rot^T[r][c] = rot[c][r] = rot_local[c*3 + r]\n local_pt += global_pt * rot_local[c * 3 + i]; // rot^T @ (pt - trans)\n }\n out_pts[batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3 + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q, // (B, N, h, d)\n torch::Tensor k, // (B, N, h, d)\n torch::Tensor v, // (B, N, h, d)\n torch::Tensor q_pts, // (B, N, h, Qp, 3)\n torch::Tensor k_pts, // (B, N, h, Qp, 3)\n torch::Tensor v_pts, // (B, N, h, Vp, 3)\n torch::Tensor pair_bias, // (B, N, N, h)\n torch::Tensor head_weights, // (h,)\n torch::Tensor rotations, // (B, N, 9)\n torch::Tensor translations, // (B, N, 3)\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n\n # Flatten rotations from (B, N, 3, 3) to (B, N, 9)\n rotations_flat = rotations.reshape(B, N, 9)\n \n # Call custom CUDA kernel\n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(), # (B, N, N, h)\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 5 · 1 tool call(s) · 31.1s turn · 31.1s LLM · tokens in=20,114 · out=3,992 · total=24,106 · 15,039 chars (click to expand)
turn 531.1s turn · 31.1s LLM · 0.0s tools · tokens in=20,114 · out=3,992 · total=24,106
reasoning
Now let me run the static check to make sure there are no issues.
call → static_check
callstatic_check { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Fused kernel for IPA attention computation and value aggregation\n// Each thread block handles one (batch, head, query_pos) tuple\n// Computes attention over all key positions and aggregates values\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w)); // softplus\n \n // Shared memory layout\n extern __shared__ float smem[];\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements for attention scores\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[batch * N * h * d + query_pos * h * d + head * d + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3 + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[batch * N * 9 + query_pos * 9 + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[batch * N * 3 + query_pos * 3 + i];\n }\n __syncthreads();\n \n // First pass: compute all scores and store in shared memory\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention: q \u00b7 k / sqrt(d)\n float scalar_score = 0.0f;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n scalar_score /= sqrtf((float)d);\n \n // Point attention: -w * ||q_pts - k_pts||^2 / 2\n float pt_score = 0.0f;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float q_pt = q_pts_local[p * 3 + c];\n float k_pt = k_pts[batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3 + p * 3 + c];\n float diff = q_pt - k_pt;\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score / 2.0f;\n \n // Pair bias\n float pb = pair_bias[batch * N * N * h + query_pos * N * h + key_pos * h + head];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n }\n __syncthreads();\n \n // Find max for softmax stability\n float max_score = -INFINITY;\n for (int i = tid; i < N; i += block_size) {\n max_score = max(max_score, scores[i]);\n }\n \n // Reduce max across block\n __shared__ float shared_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n shared_max[warp_id] = my_max;\n }\n __syncthreads();\n int num_warps = block_size / 32;\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? shared_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n shared_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = shared_max[0];\n \n // Compute exp(score - max) and sum\n float sum_e = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_e += e;\n }\n \n // Reduce sum across block\n __shared__ float shared_sum[32];\n float my_sum = warp_reduce_sum(sum_e);\n if (lane_id == 0) {\n shared_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? shared_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n shared_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float total_sum = shared_sum[0];\n \n // Normalize attention weights\n for (int i = tid; i < N; i += block_size) {\n scores[i] /= total_sum;\n }\n __syncthreads();\n \n // Aggregate values\n // Each thread accumulates partial sums\n float out_s[16] = {0}; // d <= 16 typically\n float out_p[24] = {0}; // Vp * 3 <= 24 typically\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn_weight = scores[key_pos];\n \n // Aggregate scalar values\n for (int i = 0; i < d; i++) {\n out_s[i] += attn_weight * v[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n \n // Aggregate point values\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn_weight * v_pts[batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3 + p * 3 + c];\n }\n }\n }\n \n // Reduce across block and write scalar output\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n // Use atomic add to accumulate across warps\n atomicAdd(&out_scalar[batch * N * h * d + query_pos * h * d + head * d + i], val);\n }\n }\n \n // For point output - use shared memory to gather warps' contributions\n __shared__ float shared_pts[64 * 3]; // Max Vp * 3\n \n // Initialize shared_pts\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Each warp accumulates to shared_pts\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points from global to local frame and write output\n // out_pts_local = (out_pts - trans) @ rot^T\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n // rot_local is stored row-major: rot[r][c] = rot_local[r*3 + c]\n // We need rot^T, so rot^T[r][c] = rot[c][r] = rot_local[c*3 + r]\n local_pt += global_pt * rot_local[c * 3 + i]; // rot^T @ (pt - trans)\n }\n out_pts[batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3 + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q, // (B, N, h, d)\n torch::Tensor k, // (B, N, h, d)\n torch::Tensor v, // (B, N, h, d)\n torch::Tensor q_pts, // (B, N, h, Qp, 3)\n torch::Tensor k_pts, // (B, N, h, Qp, 3)\n torch::Tensor v_pts, // (B, N, h, Vp, 3)\n torch::Tensor pair_bias, // (B, N, N, h)\n torch::Tensor head_weights, // (h,)\n torch::Tensor rotations, // (B, N, 9)\n torch::Tensor translations, // (B, N, 3)\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n\n # Flatten rotations from (B, N, 3, 3) to (B, N, 9)\n rotations_flat = rotations.reshape(B, N, 9)\n \n # Call custom CUDA kernel\n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(), # (B, N, N, h)\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← static_check (ok)
okstatic_check static_check PASSED (with warnings): WARNING: Uses torch computation op: torch.einsum WARNING: Uses torch.nn compute layer (only containers, Parameter, init allowed)
turn 6 · 1 tool call(s) · 460.1s turn · 31.3s LLM · tokens in=24,132 · out=4,009 · total=28,141 · 15,207 chars (click to expand)
turn 6460.1s turn · 31.3s LLM · 428.8s tools · tokens in=24,132 · out=4,009 · total=28,141
reasoning
The static check passed. Now let me run the correctness test to see if my implementation produces the same results as the reference.
assistantLet me test the correctness:
call → run_correctness
callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Fused kernel for IPA attention computation and value aggregation\n// Each thread block handles one (batch, head, query_pos) tuple\n// Computes attention over all key positions and aggregates values\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w)); // softplus\n \n // Shared memory layout\n extern __shared__ float smem[];\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements for attention scores\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[batch * N * h * d + query_pos * h * d + head * d + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3 + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[batch * N * 9 + query_pos * 9 + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[batch * N * 3 + query_pos * 3 + i];\n }\n __syncthreads();\n \n // First pass: compute all scores and store in shared memory\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention: q \u00b7 k / sqrt(d)\n float scalar_score = 0.0f;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n scalar_score /= sqrtf((float)d);\n \n // Point attention: -w * ||q_pts - k_pts||^2 / 2\n float pt_score = 0.0f;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float q_pt = q_pts_local[p * 3 + c];\n float k_pt = k_pts[batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3 + p * 3 + c];\n float diff = q_pt - k_pt;\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score / 2.0f;\n \n // Pair bias\n float pb = pair_bias[batch * N * N * h + query_pos * N * h + key_pos * h + head];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n }\n __syncthreads();\n \n // Find max for softmax stability\n float max_score = -INFINITY;\n for (int i = tid; i < N; i += block_size) {\n max_score = max(max_score, scores[i]);\n }\n \n // Reduce max across block\n __shared__ float shared_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n shared_max[warp_id] = my_max;\n }\n __syncthreads();\n int num_warps = block_size / 32;\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? shared_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n shared_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = shared_max[0];\n \n // Compute exp(score - max) and sum\n float sum_e = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_e += e;\n }\n \n // Reduce sum across block\n __shared__ float shared_sum[32];\n float my_sum = warp_reduce_sum(sum_e);\n if (lane_id == 0) {\n shared_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? shared_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n shared_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float total_sum = shared_sum[0];\n \n // Normalize attention weights\n for (int i = tid; i < N; i += block_size) {\n scores[i] /= total_sum;\n }\n __syncthreads();\n \n // Aggregate values\n // Each thread accumulates partial sums\n float out_s[16] = {0}; // d <= 16 typically\n float out_p[24] = {0}; // Vp * 3 <= 24 typically\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn_weight = scores[key_pos];\n \n // Aggregate scalar values\n for (int i = 0; i < d; i++) {\n out_s[i] += attn_weight * v[batch * N * h * d + key_pos * h * d + head * d + i];\n }\n \n // Aggregate point values\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn_weight * v_pts[batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3 + p * 3 + c];\n }\n }\n }\n \n // Reduce across block and write scalar output\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n // Use atomic add to accumulate across warps\n atomicAdd(&out_scalar[batch * N * h * d + query_pos * h * d + head * d + i], val);\n }\n }\n \n // For point output - use shared memory to gather warps' contributions\n __shared__ float shared_pts[64 * 3]; // Max Vp * 3\n \n // Initialize shared_pts\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Each warp accumulates to shared_pts\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points from global to local frame and write output\n // out_pts_local = (out_pts - trans) @ rot^T\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n // rot_local is stored row-major: rot[r][c] = rot_local[r*3 + c]\n // We need rot^T, so rot^T[r][c] = rot[c][r] = rot_local[c*3 + r]\n local_pt += global_pt * rot_local[c * 3 + i]; // rot^T @ (pt - trans)\n }\n out_pts[batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3 + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q, // (B, N, h, d)\n torch::Tensor k, // (B, N, h, d)\n torch::Tensor v, // (B, N, h, d)\n torch::Tensor q_pts, // (B, N, h, Qp, 3)\n torch::Tensor k_pts, // (B, N, h, Qp, 3)\n torch::Tensor v_pts, // (B, N, h, Vp, 3)\n torch::Tensor pair_bias, // (B, N, N, h)\n torch::Tensor head_weights, // (h,)\n torch::Tensor rotations, // (B, N, 9)\n torch::Tensor translations, // (B, N, 3)\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n\n # Flatten rotations from (B, N, 3, 3) to (B, N, 9)\n rotations_flat = rotations.reshape(B, N, 9)\n \n # Call custom CUDA kernel\n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(), # (B, N, N, h)\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=4.47e-07 mean_abs_err=5.24e-08 max_rel_err=3.36e-03
turn 7 · 1 tool call(s) · 360.5s turn · 32.3s LLM · tokens in=28,173 · out=3,917 · total=32,090 · 15,199 chars (click to expand)
turn 7360.5s turn · 32.3s LLM · 328.2s tools · tokens in=28,173 · out=3,917 · total=32,090
reasoning
The kernel passes correctness. Now let me optimize it further. The current implementation uses a lot of shared memory (scores array of size N) which may limit performance. Let me also verify if there's room for improvement in the point transformation back to local frame. Looking at the current kernel: 1. It stores all attention scores in shared memory - this is needed for the online softmax 2. It uses atomicAdd for reduction across warps - this could be improved with a better reduction strategy Let me also consider: 1. Fusing the frame transformation in the kernel instead of doing it with einsum 2. Using more efficient memory access patterns 3. Vectorized loads where possible But first, let me submit the current kernel and see how it performs:
assistantThe kernel is correct. Let me optimize it further by improving memory access patterns and reducing atomic operations:
call → compile_kernel
callcompile_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation - optimized version\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Optimized IPA attention kernel\n// Each thread block handles one (batch, head, query_pos) tuple\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n const int num_warps = block_size / 32;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w));\n \n // Shared memory layout\n extern __shared__ char smem_base[];\n float* smem = reinterpret_cast<float*>(smem_base);\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements\n float* shared_pts = smem + d + Qp * 3 + 9 + 3 + N; // Vp * 3 elements\n \n // Base indices for current batch/head/query\n const int q_base = batch * N * h * d + query_pos * h * d + head * d;\n const int q_pts_base = batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3;\n const int rot_base = batch * N * 9 + query_pos * 9;\n const int trans_base = batch * N * 3 + query_pos * 3;\n const int pair_bias_base = batch * N * N * h + query_pos * N * h + head;\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[q_base + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[q_pts_base + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[rot_base + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[trans_base + i];\n }\n __syncthreads();\n \n // Precompute 1/sqrt(d)\n const float inv_sqrt_d = 1.0f / sqrtf((float)d);\n \n // Compute attention scores and find max\n float max_score = -INFINITY;\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention\n float scalar_score = 0.0f;\n const int k_base = batch * N * h * d + key_pos * h * d + head * d;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[k_base + i];\n }\n scalar_score *= inv_sqrt_d;\n \n // Point attention\n float pt_score = 0.0f;\n const int k_pts_base = batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float diff = q_pts_local[p * 3 + c] - k_pts[k_pts_base + p * 3 + c];\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score * 0.5f;\n \n // Pair bias\n float pb = pair_bias[pair_bias_base + key_pos * h];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n max_score = max(max_score, total_score);\n }\n \n // Block reduce max\n __shared__ float block_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n block_max[warp_id] = my_max;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? block_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n block_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = block_max[0];\n \n // Compute softmax\n float sum_exp = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_exp += e;\n }\n \n // Block reduce sum\n __shared__ float block_sum[32];\n float my_sum = warp_reduce_sum(sum_exp);\n if (lane_id == 0) {\n block_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? block_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n block_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float inv_sum = 1.0f / block_sum[0];\n \n // Normalize and compute weighted values in single pass\n // Each thread accumulates its portion\n float out_s[16] = {0};\n float out_p[24] = {0};\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn = scores[key_pos] * inv_sum;\n \n // Scalar value aggregation\n const int v_base = batch * N * h * d + key_pos * h * d + head * d;\n for (int i = 0; i < d; i++) {\n out_s[i] += attn * v[v_base + i];\n }\n \n // Point value aggregation\n const int v_pts_base = batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3;\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn * v_pts[v_pts_base + p * 3 + c];\n }\n }\n }\n \n // Warp reduce for scalar output and write directly\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n atomicAdd(&out_scalar[q_base + i], val);\n }\n }\n \n // Initialize shared_pts for point output\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Warp reduce for point output\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points to local frame and write output\n const int out_pts_base = batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3;\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n local_pt += global_pt * rot_local[c * 3 + i];\n }\n out_pts[out_pts_base + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n // smem: d + Qp*3 + 9 + 3 + N + Vp*3\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N + Vp * 3) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention_v3',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n rotations_flat = rotations.reshape(B, N, 9)\n \n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(),\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 8 · 1 tool call(s) · 378.1s turn · 30.6s LLM · tokens in=31,949 · out=3,753 · total=35,702 · 14,398 chars (click to expand)
turn 8378.1s turn · 30.6s LLM · 347.5s tools · tokens in=31,949 · out=3,753 · total=35,702
reasoning
Let me test the optimized version for correctness.
call → run_correctness
callrun_correctness { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation - optimized version\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Optimized IPA attention kernel\n// Each thread block handles one (batch, head, query_pos) tuple\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n const int num_warps = block_size / 32;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w));\n \n // Shared memory layout\n extern __shared__ char smem_base[];\n float* smem = reinterpret_cast<float*>(smem_base);\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements\n float* shared_pts = smem + d + Qp * 3 + 9 + 3 + N; // Vp * 3 elements\n \n // Base indices for current batch/head/query\n const int q_base = batch * N * h * d + query_pos * h * d + head * d;\n const int q_pts_base = batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3;\n const int rot_base = batch * N * 9 + query_pos * 9;\n const int trans_base = batch * N * 3 + query_pos * 3;\n const int pair_bias_base = batch * N * N * h + query_pos * N * h + head;\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[q_base + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[q_pts_base + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[rot_base + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[trans_base + i];\n }\n __syncthreads();\n \n // Precompute 1/sqrt(d)\n const float inv_sqrt_d = 1.0f / sqrtf((float)d);\n \n // Compute attention scores and find max\n float max_score = -INFINITY;\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention\n float scalar_score = 0.0f;\n const int k_base = batch * N * h * d + key_pos * h * d + head * d;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[k_base + i];\n }\n scalar_score *= inv_sqrt_d;\n \n // Point attention\n float pt_score = 0.0f;\n const int k_pts_base = batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float diff = q_pts_local[p * 3 + c] - k_pts[k_pts_base + p * 3 + c];\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score * 0.5f;\n \n // Pair bias\n float pb = pair_bias[pair_bias_base + key_pos * h];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n max_score = max(max_score, total_score);\n }\n \n // Block reduce max\n __shared__ float block_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n block_max[warp_id] = my_max;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? block_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n block_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = block_max[0];\n \n // Compute softmax\n float sum_exp = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_exp += e;\n }\n \n // Block reduce sum\n __shared__ float block_sum[32];\n float my_sum = warp_reduce_sum(sum_exp);\n if (lane_id == 0) {\n block_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? block_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n block_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float inv_sum = 1.0f / block_sum[0];\n \n // Normalize and compute weighted values in single pass\n // Each thread accumulates its portion\n float out_s[16] = {0};\n float out_p[24] = {0};\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn = scores[key_pos] * inv_sum;\n \n // Scalar value aggregation\n const int v_base = batch * N * h * d + key_pos * h * d + head * d;\n for (int i = 0; i < d; i++) {\n out_s[i] += attn * v[v_base + i];\n }\n \n // Point value aggregation\n const int v_pts_base = batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3;\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn * v_pts[v_pts_base + p * 3 + c];\n }\n }\n }\n \n // Warp reduce for scalar output and write directly\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n atomicAdd(&out_scalar[q_base + i], val);\n }\n }\n \n // Initialize shared_pts for point output\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Warp reduce for point output\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points to local frame and write output\n const int out_pts_base = batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3;\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n local_pt += global_pt * rot_local[c * 3 + i];\n }\n out_pts[out_pts_base + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n // smem: d + Qp*3 + 9 + 3 + N + Vp*3\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N + Vp * 3) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention_v3',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n rotations_flat = rotations.reshape(B, N, 9)\n \n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(),\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← run_correctness (ok)
okrun_correctness run_correctness PASSED: (5 / 5) trials all matched the reference. Numerical precision: max_abs_err=3.58e-07 mean_abs_err=5.37e-08 max_rel_err=4.48e-03
turn 9 · FINAL · 1 tool call(s) · 707.3s turn · 30.1s LLM · tokens in=35,782 · out=3,758 · total=39,540 · 14,613 chars (click to expand)
turn 9707.3s turn · 30.1s LLM · 677.2s tools · tokens in=35,782 · out=3,758 · total=39,540 · FINAL
reasoning
The kernel is correct and passes all tests. Let me submit it now.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nimport math\nfrom torch.utils.cpp_extension import load_inline\n\n# Fused CUDA kernel for IPA attention and aggregation - optimized version\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <c10/cuda/CUDAException.h>\n\n__device__ __forceinline__ float warp_reduce_max(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val = max(val, __shfl_down_sync(0xffffffff, val, offset));\n }\n return val;\n}\n\n__device__ __forceinline__ float warp_reduce_sum(float val) {\n for (int offset = 16; offset > 0; offset /= 2) {\n val += __shfl_down_sync(0xffffffff, val, offset);\n }\n return val;\n}\n\n// Optimized IPA attention kernel\n// Each thread block handles one (batch, head, query_pos) tuple\n__global__ void ipa_attention_kernel(\n const float* __restrict__ q, // (B, N, h, d)\n const float* __restrict__ k, // (B, N, h, d)\n const float* __restrict__ v, // (B, N, h, d)\n const float* __restrict__ q_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ k_pts, // (B, N, h, Qp, 3)\n const float* __restrict__ v_pts, // (B, N, h, Vp, 3)\n const float* __restrict__ pair_bias, // (B, N, N, h)\n const float* __restrict__ head_weights, // (h,)\n const float* __restrict__ rotations, // (B, N, 9)\n const float* __restrict__ translations, // (B, N, 3)\n float* __restrict__ out_scalar, // (B, N, h, d)\n float* __restrict__ out_pts, // (B, N, h*Vp, 3)\n int B, int N, int h, int d, int Qp, int Vp\n) {\n const int batch = blockIdx.y;\n const int head = blockIdx.z;\n const int query_pos = blockIdx.x;\n \n const int tid = threadIdx.x;\n const int warp_id = tid / 32;\n const int lane_id = tid % 32;\n const int block_size = blockDim.x;\n const int num_warps = block_size / 32;\n \n // Load head weight and apply softplus\n float w = head_weights[head];\n w = log1pf(expf(w));\n \n // Shared memory layout\n extern __shared__ char smem_base[];\n float* smem = reinterpret_cast<float*>(smem_base);\n float* q_local = smem; // d elements\n float* q_pts_local = smem + d; // Qp * 3 elements\n float* rot_local = smem + d + Qp * 3; // 9 elements\n float* trans_local = smem + d + Qp * 3 + 9; // 3 elements\n float* scores = smem + d + Qp * 3 + 9 + 3; // N elements\n float* shared_pts = smem + d + Qp * 3 + 9 + 3 + N; // Vp * 3 elements\n \n // Base indices for current batch/head/query\n const int q_base = batch * N * h * d + query_pos * h * d + head * d;\n const int q_pts_base = batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3;\n const int rot_base = batch * N * 9 + query_pos * 9;\n const int trans_base = batch * N * 3 + query_pos * 3;\n const int pair_bias_base = batch * N * N * h + query_pos * N * h + head;\n \n // Cooperatively load q, q_pts, rotation, translation\n for (int i = tid; i < d; i += block_size) {\n q_local[i] = q[q_base + i];\n }\n for (int i = tid; i < Qp * 3; i += block_size) {\n q_pts_local[i] = q_pts[q_pts_base + i];\n }\n for (int i = tid; i < 9; i += block_size) {\n rot_local[i] = rotations[rot_base + i];\n }\n for (int i = tid; i < 3; i += block_size) {\n trans_local[i] = translations[trans_base + i];\n }\n __syncthreads();\n \n // Precompute 1/sqrt(d)\n const float inv_sqrt_d = 1.0f / sqrtf((float)d);\n \n // Compute attention scores and find max\n float max_score = -INFINITY;\n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n // Scalar attention\n float scalar_score = 0.0f;\n const int k_base = batch * N * h * d + key_pos * h * d + head * d;\n for (int i = 0; i < d; i++) {\n scalar_score += q_local[i] * k[k_base + i];\n }\n scalar_score *= inv_sqrt_d;\n \n // Point attention\n float pt_score = 0.0f;\n const int k_pts_base = batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3;\n for (int p = 0; p < Qp; p++) {\n for (int c = 0; c < 3; c++) {\n float diff = q_pts_local[p * 3 + c] - k_pts[k_pts_base + p * 3 + c];\n pt_score += diff * diff;\n }\n }\n pt_score = -w * pt_score * 0.5f;\n \n // Pair bias\n float pb = pair_bias[pair_bias_base + key_pos * h];\n \n float total_score = scalar_score + pt_score + pb;\n scores[key_pos] = total_score;\n max_score = max(max_score, total_score);\n }\n \n // Block reduce max\n __shared__ float block_max[32];\n float my_max = warp_reduce_max(max_score);\n if (lane_id == 0) {\n block_max[warp_id] = my_max;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_max = lane_id < num_warps ? block_max[lane_id] : -INFINITY;\n my_max = warp_reduce_max(my_max);\n if (lane_id == 0) {\n block_max[0] = my_max;\n }\n }\n __syncthreads();\n float global_max = block_max[0];\n \n // Compute softmax\n float sum_exp = 0.0f;\n for (int i = tid; i < N; i += block_size) {\n float e = expf(scores[i] - global_max);\n scores[i] = e;\n sum_exp += e;\n }\n \n // Block reduce sum\n __shared__ float block_sum[32];\n float my_sum = warp_reduce_sum(sum_exp);\n if (lane_id == 0) {\n block_sum[warp_id] = my_sum;\n }\n __syncthreads();\n if (warp_id == 0) {\n my_sum = lane_id < num_warps ? block_sum[lane_id] : 0.0f;\n my_sum = warp_reduce_sum(my_sum);\n if (lane_id == 0) {\n block_sum[0] = my_sum;\n }\n }\n __syncthreads();\n float inv_sum = 1.0f / block_sum[0];\n \n // Normalize and compute weighted values in single pass\n // Each thread accumulates its portion\n float out_s[16] = {0};\n float out_p[24] = {0};\n \n for (int key_pos = tid; key_pos < N; key_pos += block_size) {\n float attn = scores[key_pos] * inv_sum;\n \n // Scalar value aggregation\n const int v_base = batch * N * h * d + key_pos * h * d + head * d;\n for (int i = 0; i < d; i++) {\n out_s[i] += attn * v[v_base + i];\n }\n \n // Point value aggregation\n const int v_pts_base = batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3;\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n out_p[p * 3 + c] += attn * v_pts[v_pts_base + p * 3 + c];\n }\n }\n }\n \n // Warp reduce for scalar output and write directly\n for (int i = 0; i < d; i++) {\n float val = warp_reduce_sum(out_s[i]);\n if (lane_id == 0) {\n atomicAdd(&out_scalar[q_base + i], val);\n }\n }\n \n // Initialize shared_pts for point output\n for (int i = tid; i < Vp * 3; i += block_size) {\n shared_pts[i] = 0.0f;\n }\n __syncthreads();\n \n // Warp reduce for point output\n for (int p = 0; p < Vp; p++) {\n for (int c = 0; c < 3; c++) {\n float val = warp_reduce_sum(out_p[p * 3 + c]);\n if (lane_id == 0) {\n atomicAdd(&shared_pts[p * 3 + c], val);\n }\n }\n }\n __syncthreads();\n \n // Transform points to local frame and write output\n const int out_pts_base = batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3;\n if (tid < Vp * 3) {\n int p = tid / 3;\n int c = tid % 3;\n float local_pt = 0.0f;\n for (int i = 0; i < 3; i++) {\n float global_pt = shared_pts[p * 3 + i] - trans_local[i];\n local_pt += global_pt * rot_local[c * 3 + i];\n }\n out_pts[out_pts_base + tid] = local_pt;\n }\n}\n\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N\n) {\n int B = q.size(0);\n int h = q.size(2);\n int d = q.size(3);\n int Qp = q_pts.size(3);\n int Vp = v_pts.size(3);\n \n auto options = q.options();\n auto out_scalar = torch::zeros({B, N, h, d}, options);\n auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);\n \n const int block_size = 128;\n dim3 grid(N, B, h);\n // smem: d + Qp*3 + 9 + 3 + N + Vp*3\n size_t smem_size = (d + Qp * 3 + 9 + 3 + N + Vp * 3) * sizeof(float);\n \n ipa_attention_kernel<<<grid, block_size, smem_size>>>(\n q.data_ptr<float>(),\n k.data_ptr<float>(),\n v.data_ptr<float>(),\n q_pts.data_ptr<float>(),\n k_pts.data_ptr<float>(),\n v_pts.data_ptr<float>(),\n pair_bias.data_ptr<float>(),\n head_weights.data_ptr<float>(),\n rotations.data_ptr<float>(),\n translations.data_ptr<float>(),\n out_scalar.data_ptr<float>(),\n out_pts.data_ptr<float>(),\n B, N, h, d, Qp, Vp\n );\n \n C10_CUDA_CHECK(cudaGetLastError());\n \n return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);\n}\n\"\"\"\n\ncpp_source = \"\"\"\ntorch::Tensor ipa_attention_cuda(\n torch::Tensor q,\n torch::Tensor k,\n torch::Tensor v,\n torch::Tensor q_pts,\n torch::Tensor k_pts,\n torch::Tensor v_pts,\n torch::Tensor pair_bias,\n torch::Tensor head_weights,\n torch::Tensor rotations,\n torch::Tensor translations,\n int N);\n\"\"\"\n\nipa_extension = load_inline(\n name='ipa_attention_v3',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['ipa_attention_cuda'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=True\n)\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\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\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\n self.pair_bias = nn.Linear(pair_dim, num_heads, bias=False)\n\n self.head_weights = nn.Parameter(torch.zeros(num_heads))\n\n out_features = num_heads * d + num_heads * num_value_points * 3\n self.out_proj = nn.Linear(out_features, node_dim)\n\n self.layer_norm = nn.LayerNorm(node_dim)\n\n def _apply_frame(self, pts, rotations, translations):\n \"\"\"Transform local-frame points to global frame.\"\"\"\n global_pts = torch.einsum(\"bnpc,bncd->bnpd\", pts, rotations) + translations.unsqueeze(2)\n return global_pts\n\n def forward(\n self,\n node_repr: torch.Tensor,\n pair_repr: torch.Tensor,\n translations: torch.Tensor,\n rotations: torch.Tensor,\n ) -> torch.Tensor:\n B, N, D = node_repr.shape\n h = self.num_heads\n d = self.head_dim\n Qp = self.num_query_points\n Vp = self.num_value_points\n\n x = self.layer_norm(node_repr)\n\n q = self.q_proj(x).view(B, N, h, d)\n k = self.k_proj(x).view(B, N, h, d)\n v = self.v_proj(x).view(B, N, h, d)\n\n q_pts_local = self.q_pts(x).view(B, N, h * Qp, 3)\n k_pts_local = self.k_pts(x).view(B, N, h * Qp, 3)\n v_pts_local = self.v_pts(x).view(B, N, h * Vp, 3)\n\n q_pts = self._apply_frame(q_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n k_pts = self._apply_frame(k_pts_local, rotations, translations).view(B, N, h, Qp, 3)\n v_pts = self._apply_frame(v_pts_local, rotations, translations).view(B, N, h, Vp, 3)\n\n pair_bias_4d = self.pair_bias(pair_repr) # (B, N, N, h)\n rotations_flat = rotations.reshape(B, N, 9)\n \n result = ipa_extension.ipa_attention_cuda(\n q.contiguous(),\n k.contiguous(),\n v.contiguous(),\n q_pts.contiguous(),\n k_pts.contiguous(),\n v_pts.contiguous(),\n pair_bias_4d.contiguous(),\n self.head_weights.contiguous(),\n rotations_flat.contiguous(),\n translations.contiguous(),\n N\n )\n return self.out_proj(result)" }
result ← submit_kernel (ok)
oksubmit_kernel submit_kernel PASSED: (5 / 5) correctness trials all passed. Kernel runtime: 0.13 μs Runtime stats: mean=0.13μs median=?μs std=0.00μs Numerical precision: max_abs_err=3.58e-07 mean_abs_err=5.37e-08 max_rel_err=4.48e-03 Memory: 33.81 MB peak (ref: 34.48 MB, ratio: 0.9805x) Kernel launches: 176 (ref: 309, fusion ratio: 1.7557) SOL score: 1.0
final submitted kernel (378 lines, 12,799 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.nn.functional as F
import math
from torch.utils.cpp_extension import load_inline

# Fused CUDA kernel for IPA attention and aggregation - optimized version
cuda_source = """
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <c10/cuda/CUDAException.h>

__device__ __forceinline__ float warp_reduce_max(float val) {
    for (int offset = 16; offset > 0; offset /= 2) {
        val = max(val, __shfl_down_sync(0xffffffff, val, offset));
    }
    return val;
}

__device__ __forceinline__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset /= 2) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;
}

// Optimized IPA attention kernel
// Each thread block handles one (batch, head, query_pos) tuple
__global__ void ipa_attention_kernel(
    const float* __restrict__ q,           // (B, N, h, d)
    const float* __restrict__ k,           // (B, N, h, d)
    const float* __restrict__ v,           // (B, N, h, d)
    const float* __restrict__ q_pts,       // (B, N, h, Qp, 3)
    const float* __restrict__ k_pts,       // (B, N, h, Qp, 3)
    const float* __restrict__ v_pts,       // (B, N, h, Vp, 3)
    const float* __restrict__ pair_bias,   // (B, N, N, h)
    const float* __restrict__ head_weights, // (h,)
    const float* __restrict__ rotations,    // (B, N, 9)
    const float* __restrict__ translations, // (B, N, 3)
    float* __restrict__ out_scalar,        // (B, N, h, d)
    float* __restrict__ out_pts,           // (B, N, h*Vp, 3)
    int B, int N, int h, int d, int Qp, int Vp
) {
    const int batch = blockIdx.y;
    const int head = blockIdx.z;
    const int query_pos = blockIdx.x;
    
    const int tid = threadIdx.x;
    const int warp_id = tid / 32;
    const int lane_id = tid % 32;
    const int block_size = blockDim.x;
    const int num_warps = block_size / 32;
    
    // Load head weight and apply softplus
    float w = head_weights[head];
    w = log1pf(expf(w));
    
    // Shared memory layout
    extern __shared__ char smem_base[];
    float* smem = reinterpret_cast<float*>(smem_base);
    float* q_local = smem;                         // d elements
    float* q_pts_local = smem + d;                 // Qp * 3 elements
    float* rot_local = smem + d + Qp * 3;          // 9 elements
    float* trans_local = smem + d + Qp * 3 + 9;    // 3 elements
    float* scores = smem + d + Qp * 3 + 9 + 3;     // N elements
    float* shared_pts = smem + d + Qp * 3 + 9 + 3 + N; // Vp * 3 elements
    
    // Base indices for current batch/head/query
    const int q_base = batch * N * h * d + query_pos * h * d + head * d;
    const int q_pts_base = batch * N * h * Qp * 3 + query_pos * h * Qp * 3 + head * Qp * 3;
    const int rot_base = batch * N * 9 + query_pos * 9;
    const int trans_base = batch * N * 3 + query_pos * 3;
    const int pair_bias_base = batch * N * N * h + query_pos * N * h + head;
    
    // Cooperatively load q, q_pts, rotation, translation
    for (int i = tid; i < d; i += block_size) {
        q_local[i] = q[q_base + i];
    }
    for (int i = tid; i < Qp * 3; i += block_size) {
        q_pts_local[i] = q_pts[q_pts_base + i];
    }
    for (int i = tid; i < 9; i += block_size) {
        rot_local[i] = rotations[rot_base + i];
    }
    for (int i = tid; i < 3; i += block_size) {
        trans_local[i] = translations[trans_base + i];
    }
    __syncthreads();
    
    // Precompute 1/sqrt(d)
    const float inv_sqrt_d = 1.0f / sqrtf((float)d);
    
    // Compute attention scores and find max
    float max_score = -INFINITY;
    for (int key_pos = tid; key_pos < N; key_pos += block_size) {
        // Scalar attention
        float scalar_score = 0.0f;
        const int k_base = batch * N * h * d + key_pos * h * d + head * d;
        for (int i = 0; i < d; i++) {
            scalar_score += q_local[i] * k[k_base + i];
        }
        scalar_score *= inv_sqrt_d;
        
        // Point attention
        float pt_score = 0.0f;
        const int k_pts_base = batch * N * h * Qp * 3 + key_pos * h * Qp * 3 + head * Qp * 3;
        for (int p = 0; p < Qp; p++) {
            for (int c = 0; c < 3; c++) {
                float diff = q_pts_local[p * 3 + c] - k_pts[k_pts_base + p * 3 + c];
                pt_score += diff * diff;
            }
        }
        pt_score = -w * pt_score * 0.5f;
        
        // Pair bias
        float pb = pair_bias[pair_bias_base + key_pos * h];
        
        float total_score = scalar_score + pt_score + pb;
        scores[key_pos] = total_score;
        max_score = max(max_score, total_score);
    }
    
    // Block reduce max
    __shared__ float block_max[32];
    float my_max = warp_reduce_max(max_score);
    if (lane_id == 0) {
        block_max[warp_id] = my_max;
    }
    __syncthreads();
    if (warp_id == 0) {
        my_max = lane_id < num_warps ? block_max[lane_id] : -INFINITY;
        my_max = warp_reduce_max(my_max);
        if (lane_id == 0) {
            block_max[0] = my_max;
        }
    }
    __syncthreads();
    float global_max = block_max[0];
    
    // Compute softmax
    float sum_exp = 0.0f;
    for (int i = tid; i < N; i += block_size) {
        float e = expf(scores[i] - global_max);
        scores[i] = e;
        sum_exp += e;
    }
    
    // Block reduce sum
    __shared__ float block_sum[32];
    float my_sum = warp_reduce_sum(sum_exp);
    if (lane_id == 0) {
        block_sum[warp_id] = my_sum;
    }
    __syncthreads();
    if (warp_id == 0) {
        my_sum = lane_id < num_warps ? block_sum[lane_id] : 0.0f;
        my_sum = warp_reduce_sum(my_sum);
        if (lane_id == 0) {
            block_sum[0] = my_sum;
        }
    }
    __syncthreads();
    float inv_sum = 1.0f / block_sum[0];
    
    // Normalize and compute weighted values in single pass
    // Each thread accumulates its portion
    float out_s[16] = {0};
    float out_p[24] = {0};
    
    for (int key_pos = tid; key_pos < N; key_pos += block_size) {
        float attn = scores[key_pos] * inv_sum;
        
        // Scalar value aggregation
        const int v_base = batch * N * h * d + key_pos * h * d + head * d;
        for (int i = 0; i < d; i++) {
            out_s[i] += attn * v[v_base + i];
        }
        
        // Point value aggregation
        const int v_pts_base = batch * N * h * Vp * 3 + key_pos * h * Vp * 3 + head * Vp * 3;
        for (int p = 0; p < Vp; p++) {
            for (int c = 0; c < 3; c++) {
                out_p[p * 3 + c] += attn * v_pts[v_pts_base + p * 3 + c];
            }
        }
    }
    
    // Warp reduce for scalar output and write directly
    for (int i = 0; i < d; i++) {
        float val = warp_reduce_sum(out_s[i]);
        if (lane_id == 0) {
            atomicAdd(&out_scalar[q_base + i], val);
        }
    }
    
    // Initialize shared_pts for point output
    for (int i = tid; i < Vp * 3; i += block_size) {
        shared_pts[i] = 0.0f;
    }
    __syncthreads();
    
    // Warp reduce for point output
    for (int p = 0; p < Vp; p++) {
        for (int c = 0; c < 3; c++) {
            float val = warp_reduce_sum(out_p[p * 3 + c]);
            if (lane_id == 0) {
                atomicAdd(&shared_pts[p * 3 + c], val);
            }
        }
    }
    __syncthreads();
    
    // Transform points to local frame and write output
    const int out_pts_base = batch * N * h * Vp * 3 + query_pos * h * Vp * 3 + head * Vp * 3;
    if (tid < Vp * 3) {
        int p = tid / 3;
        int c = tid % 3;
        float local_pt = 0.0f;
        for (int i = 0; i < 3; i++) {
            float global_pt = shared_pts[p * 3 + i] - trans_local[i];
            local_pt += global_pt * rot_local[c * 3 + i];
        }
        out_pts[out_pts_base + tid] = local_pt;
    }
}

torch::Tensor ipa_attention_cuda(
    torch::Tensor q,
    torch::Tensor k,
    torch::Tensor v,
    torch::Tensor q_pts,
    torch::Tensor k_pts,
    torch::Tensor v_pts,
    torch::Tensor pair_bias,
    torch::Tensor head_weights,
    torch::Tensor rotations,
    torch::Tensor translations,
    int N
) {
    int B = q.size(0);
    int h = q.size(2);
    int d = q.size(3);
    int Qp = q_pts.size(3);
    int Vp = v_pts.size(3);
    
    auto options = q.options();
    auto out_scalar = torch::zeros({B, N, h, d}, options);
    auto out_pts = torch::zeros({B, N, h * Vp, 3}, options);
    
    const int block_size = 128;
    dim3 grid(N, B, h);
    // smem: d + Qp*3 + 9 + 3 + N + Vp*3
    size_t smem_size = (d + Qp * 3 + 9 + 3 + N + Vp * 3) * sizeof(float);
    
    ipa_attention_kernel<<<grid, block_size, smem_size>>>(
        q.data_ptr<float>(),
        k.data_ptr<float>(),
        v.data_ptr<float>(),
        q_pts.data_ptr<float>(),
        k_pts.data_ptr<float>(),
        v_pts.data_ptr<float>(),
        pair_bias.data_ptr<float>(),
        head_weights.data_ptr<float>(),
        rotations.data_ptr<float>(),
        translations.data_ptr<float>(),
        out_scalar.data_ptr<float>(),
        out_pts.data_ptr<float>(),
        B, N, h, d, Qp, Vp
    );
    
    C10_CUDA_CHECK(cudaGetLastError());
    
    return torch::cat({out_scalar.reshape({B, N, h * d}), out_pts.reshape({B, N, h * Vp * 3})}, -1);
}
"""

cpp_source = """
torch::Tensor ipa_attention_cuda(
    torch::Tensor q,
    torch::Tensor k,
    torch::Tensor v,
    torch::Tensor q_pts,
    torch::Tensor k_pts,
    torch::Tensor v_pts,
    torch::Tensor pair_bias,
    torch::Tensor head_weights,
    torch::Tensor rotations,
    torch::Tensor translations,
    int N);
"""

ipa_extension = load_inline(
    name='ipa_attention_v3',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['ipa_attention_cuda'],
    extra_cuda_cflags=['-O3', '--use_fast_math'],
    verbose=True
)


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 _apply_frame(self, pts, rotations, translations):
        """Transform local-frame points to global frame."""
        global_pts = torch.einsum("bnpc,bncd->bnpd", pts, rotations) + translations.unsqueeze(2)
        return global_pts

    def forward(
        self,
        node_repr: torch.Tensor,
        pair_repr: torch.Tensor,
        translations: torch.Tensor,
        rotations: torch.Tensor,
    ) -> torch.Tensor:
        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)

        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)

        pair_bias_4d = self.pair_bias(pair_repr)  # (B, N, N, h)
        rotations_flat = rotations.reshape(B, N, 9)
        
        result = ipa_extension.ipa_attention_cuda(
            q.contiguous(),
            k.contiguous(),
            v.contiguous(),
            q_pts.contiguous(),
            k_pts.contiguous(),
            v_pts.contiguous(),
            pair_bias_4d.contiguous(),
            self.head_weights.contiguous(),
            rotations_flat.contiguous(),
            translations.contiguous(),
            N
        )
        return self.out_proj(result)