call → compile_kernel
callcompile_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#include <cmath>\n\n__device__ float device_sigmoid(float x) { return 1.0f / (1.0f + expf(-x)); }\n\n__global__ void triangular_update_kernel(\n const float* __restrict__ pair,\n const float* __restrict__ mask,\n const float* __restrict__ lnw, const float* __restrict__ lnb,\n const float* __restrict__ paw, const float* __restrict__ pab,\n const float* __restrict__ gaw, const float* __restrict__ gab,\n const float* __restrict__ pbw, const float* __restrict__ pbb,\n const float* __restrict__ gbw, const float* __restrict__ gbb,\n const float* __restrict__ lnoutw, const float* __restrict__ lnoutb,\n const float* __restrict__ opw, const float* __restrict__ opb,\n const float* __restrict__ ogw, const float* __restrict__ ogb,\n float* __restrict__ out,\n int B, int N, int D, int H\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * N * N * D;\n if (idx >= total) return;\n int d_out = idx % D;\n int j = (idx / D) % N;\n int i = (idx / (D * N)) % N;\n int b = idx / (D * N * N);\n \n float eps = 1e-5f;\n // LN in for (b,i,j)\n float mean = 0.0f;\n #pragma unroll 8\n for (int d=0; d<D; d++) mean += pair[((b * N + i) * N + j) * D + d];\n mean /= D;\n float var = 0.0f;\n #pragma unroll 8\n for (int d=0; d<D; d++) {\n float v = pair[((b * N + i) * N + j) * D + d] - mean;\n var += v * v;\n }\n var /= D;\n float rstd = rsqrtf(var + eps);\n \n float x_d = (pair[((b * N + i) * N + j) * D + d_out] - mean) * rstd * lnw[d_out] + lnb[d_out];\n \n // Now to compute the out, we need to project to H, aggregate, project back\n // So we must compute the full out vector for (b,i,j)\n // Since we are per d_out, we redo work, but for correctness OK\n // First compute a and b for all k, all h\n float res = 0.0f;\n for (int h=0; h<H; h++) {\n float sum = 0.0f;\n for (int k=0; k<N; k++) {\n // compute a_h for (b,i,k)\n float mean_ak = 0.0f;\n for (int d=0; d<D; d++) mean_ak += pair[((b*N + i)*N + k)*D + d];\n mean_ak /= D;\n float var_ak = 0.0f;\n for (int d=0; d<D; d++) {\n float v = pair[((b*N + i)*N + k)*D + d] - mean_ak;\n var_ak += v*v;\n }\n var_ak /= D;\n float rstd_ak = rsqrtf(var_ak + eps);\n float x_ak = 0.0f; // for proj we need all\n // This is too slow and complex - redo all projections\n }\n }\n out[idx] = pair[idx]; // still placeholder\n}\n\ntorch::Tensor triangular_update_cuda(torch::Tensor pair_repr, torch::Tensor pair_mask,\n torch::Tensor ln_in_w, torch::Tensor ln_in_b,\n torch::Tensor pa_w, torch::Tensor pa_b, torch::Tensor ga_w, torch::Tensor ga_b,\n torch::Tensor pb_w, torch::Tensor pb_b, torch::Tensor gb_w, torch::Tensor gb_b,\n torch::Tensor ln_out_w, torch::Tensor ln_out_b,\n torch::Tensor op_w, torch::Tensor op_b, torch::Tensor og_w, torch::Tensor og_b) {\n auto B = pair_repr.size(0);\n auto N = pair_repr.size(1);\n auto D = pair_repr.size(3);\n auto H = pa_w.size(0);\n auto output = torch::zeros_like(pair_repr);\n int total = B*N*N*D;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n triangular_update_kernel<<<blocks, threads>>>(\n pair_repr.data_ptr<float>(), pair_mask.data_ptr<float>(),\n ln_in_w.data_ptr<float>(), ln_in_b.data_ptr<float>(),\n pa_w.data_ptr<float>(), pa_b.data_ptr<float>(),\n ga_w.data_ptr<float>(), ga_b.data_ptr<float>(),\n pb_w.data_ptr<float>(), pb_b.data_ptr<float>(),\n gb_w.data_ptr<float>(), gb_b.data_ptr<float>(),\n ln_out_w.data_ptr<float>(), ln_out_b.data_ptr<float>(),\n op_w.data_ptr<float>(), op_b.data_ptr<float>(),\n og_w.data_ptr<float>(), og_b.data_ptr<float>(),\n output.data_ptr<float>(), B, N, D, H);\n return output;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor triangular_update_cuda(torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor);\n'''\n\nmodule = load_inline(\n name='tri_update',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['triangular_update_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm_in = nn.LayerNorm(pair_dim)\n self.proj_a = nn.Linear(pair_dim, hidden_dim)\n self.gate_a = nn.Linear(pair_dim, hidden_dim)\n self.proj_b = nn.Linear(pair_dim, hidden_dim)\n self.gate_b = nn.Linear(pair_dim, hidden_dim)\n self.layer_norm_out = nn.LayerNorm(hidden_dim)\n self.out_proj = nn.Linear(hidden_dim, pair_dim)\n self.out_gate = nn.Linear(pair_dim, pair_dim)\n self.module = module\n\n def forward(self, pair_repr, pair_mask):\n return self.module.triangular_update_cuda(\n pair_repr.contiguous(), pair_mask.contiguous(),\n self.layer_norm_in.weight, self.layer_norm_in.bias,\n self.proj_a.weight, self.proj_a.bias, self.gate_a.weight, self.gate_a.bias,\n self.proj_b.weight, self.proj_b.bias, self.gate_b.weight, self.gate_b.bias,\n self.layer_norm_out.weight, self.layer_norm_out.bias,\n self.out_proj.weight, self.out_proj.bias, self.out_gate.weight, self.out_gate.bias)\n"
}