turn 0425.0s turn · 33.2s LLM · 391.9s tools · tokens in=1,429 · out=1,585 · total=3,860 · FINAL
reasoning
The task is to replace the forward computation of the PyTorch model with a custom CUDA kernel. The model is an outer-product mean from AlphaFold2.
I need to implement a custom CUDA kernel for the outer-product mean computation in the AlphaFold2 model. The `ModelNew` class should match the original constructor and forward signature, replacing the forward computation with a CUDA kernel for speed, while ensuring equivalent outputs.
The computation involves:
- Applying layer norm to `msa (B, S, N, msa_dim)`
- Two linear projections to `hidden_dim`: `a` and `b`
- Multiplying by an unsqueezed `msa_mask`
- Performing an einsum: `outer = torch.
assistant```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void outer_mean_kernel(
const float* __restrict__ a,
const float* __restrict__ b,
const float* __restrict__ denom,
float* __restrict__ out,
int B, int S, int N, int H
) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int total = B * N * N * H * H;
if (idx >= total) return;
int h2 = idx % H;
int h1 = (idx / H) % H;
int j = (idx / (H*H)) % N;
int i = (idx / (H*H*N)) % N;
int batch = idx / (H*H*N*N);
float sum = 0.0f;
float d = denom[batch * N + i];
if (d < 1.0f) d = 1.0f;
for (int s = 0; s < S; s++) {
float aval = a[((batch * S + s) * N + i) * H + h1];
float bval = b[((batch * S + s) * N + j) * H + h2];
sum += aval * bval;
}
out[idx] = sum / d;
}
torch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask) {
auto B = a.size(0);
auto S = a.size(1);
auto N = a.size(2);
auto H = a.size(3);
auto denom = msa_mask.sum(1); // (B, N)
auto options = torch::TensorOptions().dtype(torch::kFloat32).device(a.device());
auto out = torch::zeros({B, N, N, H*H}, options);
int total = B * N * N * H * H;
int threads = 256;
int blocks = (total + threads - 1) / threads;
outer_mean_kernel<<<blocks, threads>>>(
a.data_ptr<float>(),
b.data_ptr<float>(),
denom.data_ptr<float>(),
out.data_ptr<float>(),
B, S, N, H
);
return out;
}
'''
cpp_source = '''
torch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask);
'''
module = load_inline(
name='outer_mean_ext',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['outer_mean_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, msa_dim, pair_dim, hidden_dim=32):
super().__init__()
self.layer_norm = nn.LayerNorm(msa_dim)
self.proj_a = nn.Linear(msa_dim, hidden_dim)
self.proj_b = nn.Linear(msa_dim, hidden_dim)
self.out_proj = nn.Linear(hidden_dim * hidden_dim, pair_dim)
self.ext = module
def forward(self, msa_repr: torch.Tensor, msa_mask: torch.Tensor) -> torch.Tensor:
x = self.layer_norm(msa_repr)
a = self.proj_a(x)
b = self.proj_b(x)
mask = msa_mask.unsqueeze(-1)
a = a * mask
b = b * mask
outer = self.ext.outer_mean_cuda(a, b, msa_mask)
return self.out_proj(outer)
```
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void outer_mean_kernel(\n const float* __restrict__ a,\n const float* __restrict__ b,\n const float* __restrict__ denom,\n float* __restrict__ out,\n int B, int S, int N, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * H * H;\n if (idx >= total) return;\n\n int h2 = idx % H;\n int h1 = (idx / H) % H;\n int j = (idx / (H*H)) % N;\n int i = (idx / (H*H*N)) % N;\n int batch = idx / (H*H*N*N);\n\n float sum = 0.0f;\n float d = denom[batch * N + i];\n if (d < 1.0f) d = 1.0f;\n\n for (int s = 0; s < S; s++) {\n float aval = a[((batch * S + s) * N + i) * H + h1];\n float bval = b[((batch * S + s) * N + j) * H + h2];\n sum += aval * bval;\n }\n out[idx] = sum / d;\n}\n\ntorch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask) {\n auto B = a.size(0);\n auto S = a.size(1);\n auto N = a.size(2);\n auto H = a.size(3);\n\n auto denom = msa_mask.sum(1); // (B, N)\n auto options = torch::TensorOptions().dtype(torch::kFloat32).device(a.device());\n auto out = torch::zeros({B, N, N, H*H}, options);\n\n int total = B * N * N * H * H;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n\n outer_mean_kernel<<<blocks, threads>>>(\n a.data_ptr<float>(),\n b.data_ptr<float>(),\n denom.data_ptr<float>(),\n out.data_ptr<float>(),\n B, S, N, H\n );\n\n return out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor outer_mean_cuda(torch::Tensor a, torch::Tensor b, torch::Tensor msa_mask);\n'''\n\nmodule = load_inline(\n name='outer_mean_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['outer_mean_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, msa_dim, pair_dim, hidden_dim=32):\n super().__init__()\n self.layer_norm = nn.LayerNorm(msa_dim)\n self.proj_a = nn.Linear(msa_dim, hidden_dim)\n self.proj_b = nn.Linear(msa_dim, hidden_dim)\n self.out_proj = nn.Linear(hidden_dim * hidden_dim, pair_dim)\n self.ext = module\n\n def forward(self, msa_repr: torch.Tensor, msa_mask: torch.Tensor) -> torch.Tensor:\n x = self.layer_norm(msa_repr)\n a = self.proj_a(x)\n b = self.proj_b(x)\n\n mask = msa_mask.unsqueeze(-1)\n a = a * mask\n b = b * mask\n\n outer = self.ext.outer_mean_cuda(a, b, msa_mask)\n return self.out_proj(outer)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 0.12 μs
Runtime stats: mean=0.12μs median=?μs std=0.00μs
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
Memory: 42.41 MB peak (ref: 49.91 MB, ratio: 0.8497x)
Kernel launches: 38 (ref: 75, fusion ratio: 1.9737)
Energy: 229.36 mJ/run (ref: 0.0 mJ, ratio: -1x)
SOL score: 0.0975