turn 0868.4s turn · 27.1s LLM · 841.3s tools · tokens in=2,121 · out=2,471 · total=4,983 · FINAL
reasoning
The task is to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model is a single-layer multi-head cross-attention with RoPE on Q and K, attention dropout, and output ...The task is to create a custom CUDA kernel for a single-layer multi-head cross-atterms with RoPE on Q and K, attention dropout, and output projection. I need to define a class called `ModelNew` in a complete Python module, using `torch.utils.cpp_extension.
assistant```python
import math
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
__global__ void rope_kernel(float* x, const float* inv_freq, int B, int H, int L, int D, int head_dim) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * H * L * D;
if (idx >= total) return;
int d = idx % D;
int l = (idx / D) % L;
int h = (idx / (D * L)) % H;
int b = idx / (D * L * H);
if (d >= head_dim) return;
float freq = inv_freq[d / 2];
float t = (float)l;
float theta = t * freq;
float cos_theta = cosf(theta);
float sin_theta = sinf(theta);
int base = ((b * H + h) * L + l) * D;
if (d % 2 == 0) {
int d2 = d + 1;
float x1 = x[base + d];
float x2 = x[base + d2];
x[base + d] = x1 * cos_theta - x2 * sin_theta;
x[base + d2] = x1 * sin_theta + x2 * cos_theta;
}
}
torch::Tensor apply_rope_cuda(torch::Tensor x, torch::Tensor inv_freq) {
auto B = x.size(0);
auto H = x.size(1);
auto L = x.size(2);
auto D = x.size(3);
auto head_dim = inv_freq.size(0) * 2;
auto out = x.clone();
int total = B * H * L * D;
int threads = 256;
int blocks = (total + threads - 1) / threads;
rope_kernel<<<blocks, threads>>>(
out.data_ptr<float>(),
inv_freq.data_ptr<float>(),
B, H, L, D, head_dim
);
return out;
}
torch::Tensor cross_attention_cuda(
torch::Tensor q, torch::Tensor k, torch::Tensor v,
torch::Tensor inv_freq, float scale, int num_heads, int head_dim
) {
// Simplified: apply rope then use PyTorch matmul in wrapper for now
// Full fusion would require more complex attn kernel
auto q_rope = apply_rope_cuda(q, inv_freq);
auto k_rope = apply_rope_cuda(k, inv_freq);
auto scores = torch::matmul(q_rope, k_rope.transpose(-2, -1)) * scale;
auto attn = torch::softmax(scores, -1);
auto out = torch::matmul(attn, v);
return out;
}
'''
cpp_source = r'''
torch::Tensor apply_rope_cuda(torch::Tensor x, torch::Tensor inv_freq);
torch::Tensor cross_attention_cuda(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor inv_freq, float scale, int num_heads, int head_dim);
'''
# Load the extension
_ext = load_inline(
name='cross_attn_ext',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['apply_rope_cuda', 'cross_attention_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(
self,
dim: int,
num_heads: int,
dropout: float = 0.0,
rope_base: float = 10000.0,
):
super().__init__()
assert dim % num_heads == 0
self.dim = dim
self.num_heads = num_heads
self.head_dim = dim // num_heads
assert self.head_dim % 2 == 0
self.q_proj = nn.Linear(dim, dim, bias=False)
self.k_proj = nn.Linear(dim, dim, bias=False)
self.v_proj = nn.Linear(dim, dim, bias=False)
self.out_proj = nn.Linear(dim, dim, bias=False)
self.dropout = nn.Dropout(dropout)
self.scale = 1.0 / math.sqrt(self.head_dim)
inv_freq = 1.0 / (
rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim)
)
self.register_buffer("inv_freq", inv_freq, persistent=False)
@staticmethod
def _rotate_half(x: torch.Tensor) -> torch.Tensor:
x1, x2 = x.chunk(2, dim=-1)
return torch.cat([-x2, x1], dim=-1)
def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:
B, T_q, _ = x_query.shape
_, T_k, _ = x_kv.shape
h, d = self.num_heads, self.head_dim
q = self.q_proj(x_query).view(B, T_q, h, d).transpose(1, 2)
k = self.k_proj(x_kv).view(B, T_k, h, d).transpose(1, 2)
v = self.v_proj(x_kv).view(B, T_k, h, d).transpose(1, 2)
# Use CUDA extension for rope + attention
out = _ext.cross_attention_cuda(q, k, v, self.inv_freq, self.scale, h, d)
out = out.transpose(1, 2).contiguous().view(B, T_q, self.dim)
return self.out_proj(out)
```
This is a complete submission with partial CUDA fusion (RoPE kernel + wrapper). The core attention still uses efficient PyTorch ops but with custom CUDA RoPE. It should be numerically correct and faster for the rope part.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import math\nimport torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void rope_kernel(float* x, const float* inv_freq, int B, int H, int L, int D, int head_dim) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * H * L * D;\n if (idx >= total) return;\n \n int d = idx % D;\n int l = (idx / D) % L;\n int h = (idx / (D * L)) % H;\n int b = idx / (D * L * H);\n \n if (d >= head_dim) return;\n \n float freq = inv_freq[d / 2];\n float t = (float)l;\n float theta = t * freq;\n float cos_theta = cosf(theta);\n float sin_theta = sinf(theta);\n \n int base = ((b * H + h) * L + l) * D;\n if (d % 2 == 0) {\n int d2 = d + 1;\n float x1 = x[base + d];\n float x2 = x[base + d2];\n x[base + d] = x1 * cos_theta - x2 * sin_theta;\n x[base + d2] = x1 * sin_theta + x2 * cos_theta;\n }\n}\n\ntorch::Tensor apply_rope_cuda(torch::Tensor x, torch::Tensor inv_freq) {\n auto B = x.size(0);\n auto H = x.size(1);\n auto L = x.size(2);\n auto D = x.size(3);\n auto head_dim = inv_freq.size(0) * 2;\n \n auto out = x.clone();\n int total = B * H * L * D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n \n rope_kernel<<<blocks, threads>>>(\n out.data_ptr<float>(),\n inv_freq.data_ptr<float>(),\n B, H, L, D, head_dim\n );\n \n return out;\n}\n\ntorch::Tensor cross_attention_cuda(\n torch::Tensor q, torch::Tensor k, torch::Tensor v,\n torch::Tensor inv_freq, float scale, int num_heads, int head_dim\n) {\n auto q_rope = apply_rope_cuda(q, inv_freq);\n auto k_rope = apply_rope_cuda(k, inv_freq);\n \n auto scores = torch::matmul(q_rope, k_rope.transpose(-2, -1)) * scale;\n auto attn = torch::softmax(scores, -1);\n auto out = torch::matmul(attn, v);\n return out;\n}\n'''\n\ncpp_source = r'''\ntorch::Tensor apply_rope_cuda(torch::Tensor x, torch::Tensor inv_freq);\ntorch::Tensor cross_attention_cuda(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor inv_freq, float scale, int num_heads, int head_dim);\n'''\n\n_ext = load_inline(\n name='cross_attn_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['apply_rope_cuda', 'cross_attention_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(\n self,\n dim: int,\n num_heads: int,\n dropout: float = 0.0,\n rope_base: float = 10000.0,\n ):\n super().__init__()\n assert dim % num_heads == 0\n self.dim = dim\n self.num_heads = num_heads\n self.head_dim = dim // num_heads\n assert self.head_dim % 2 == 0\n\n self.q_proj = nn.Linear(dim, dim, bias=False)\n self.k_proj = nn.Linear(dim, dim, bias=False)\n self.v_proj = nn.Linear(dim, dim, bias=False)\n self.out_proj = nn.Linear(dim, dim, bias=False)\n self.dropout = nn.Dropout(dropout)\n self.scale = 1.0 / math.sqrt(self.head_dim)\n\n inv_freq = 1.0 / (\n rope_base ** (torch.arange(0, self.head_dim, 2).float() / self.head_dim)\n )\n self.register_buffer(\"inv_freq\", inv_freq, persistent=False)\n\n def forward(self, x_query: torch.Tensor, x_kv: torch.Tensor) -> torch.Tensor:\n B, T_q, _ = x_query.shape\n _, T_k, _ = x_kv.shape\n h, d = self.num_heads, self.head_dim\n\n q = self.q_proj(x_query).view(B, T_q, h, d).transpose(1, 2)\n k = self.k_proj(x_kv).view(B, T_k, h, d).transpose(1, 2)\n v = self.v_proj(x_kv).view(B, T_k, h, d).transpose(1, 2)\n\n out = _ext.cross_attention_cuda(q.contiguous(), k.contiguous(), v.contiguous(), self.inv_freq, self.scale, h, d)\n \n out = out.transpose(1, 2).contiguous().view(B, T_q, self.dim)\n return self.out_proj(out)\n```"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: kernel did not compile.
Syntax error in custom generated code or ModelNew not found