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.
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"
}