call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncpp_source = r\"\"\"\n#include <torch/extension.h>\ntorch::Tensor contact_forward(torch::Tensor pair_repr, torch::Tensor ln_weight, torch::Tensor ln_bias,\n torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3);\n\"\"\"\n\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n#define C_DIM 64\n#define H_DIM 32\n#define WARPS 4\n#define BLOCK 128\n__device__ __forceinline__ float warp_sum(float v){\n v += __shfl_down_sync(0xffffffff, v, 16);\n v += __shfl_down_sync(0xffffffff, v, 8);\n v += __shfl_down_sync(0xffffffff, v, 4);\n v += __shfl_down_sync(0xffffffff, v, 2);\n v += __shfl_down_sync(0xffffffff, v, 1);\n return v;\n}\n__global__ __launch_bounds__(BLOCK) void contact_kernel_tri(\n const float* __restrict__ inp, const float* __restrict__ ln_w, const float* __restrict__ ln_b,\n const float* __restrict__ w1, const float* __restrict__ bias1, const float* __restrict__ w2,\n const float* __restrict__ bias2, const float* __restrict__ w3, const float* __restrict__ bias3,\n float* __restrict__ out, int N, int tri, int total_tri)\n{\n __shared__ float sx[WARPS*C_DIM];\n __shared__ float h1[WARPS*H_DIM];\n __shared__ float h2[WARPS*H_DIM];\n int lane = threadIdx.x & 31;\n int warp = threadIdx.x >> 5;\n int wid = blockIdx.x * WARPS + warp;\n if (wid >= total_tri) return;\n int b = wid / tri;\n int kk = wid - b * tri;\n float nf = (float)N;\n float val = (2.0f*nf + 1.0f);\n int i = (int)floorf((val - sqrtf(val*val - 8.0f*(float)kk)) * 0.5f);\n int start = i * N - (i * (i - 1)) / 2;\n int j = i + (kk - start);\n\n long row1 = (((long)b * N + i) * N + j) * C_DIM;\n long row2 = (((long)b * N + j) * N + i) * C_DIM;\n float a0=inp[row1+lane], a1=inp[row1+lane+32];\n float c0=inp[row2+lane], c1=inp[row2+lane+32];\n float mean1=warp_sum(a0+a1)*(1.0f/C_DIM);\n float mean2=warp_sum(c0+c1)*(1.0f/C_DIM);\n mean1=__shfl_sync(0xffffffff, mean1, 0);\n mean2=__shfl_sync(0xffffffff, mean2, 0);\n float d0=a0-mean1, d1=a1-mean1, e0=c0-mean2, e1=c1-mean2;\n float rstd1=rsqrtf(warp_sum(d0*d0+d1*d1)*(1.0f/C_DIM)+1.0e-5f);\n float rstd2=rsqrtf(warp_sum(e0*e0+e1*e1)*(1.0f/C_DIM)+1.0e-5f);\n rstd1=__shfl_sync(0xffffffff, rstd1, 0);\n rstd2=__shfl_sync(0xffffffff, rstd2, 0);\n float* s=sx+warp*C_DIM;\n s[lane] = 0.5f*(d0*rstd1*ln_w[lane]+ln_b[lane] + e0*rstd2*ln_w[lane]+ln_b[lane]);\n s[lane+32] = 0.5f*(d1*rstd1*ln_w[lane+32]+ln_b[lane+32] + e1*rstd2*ln_w[lane+32]+ln_b[lane+32]);\n __syncwarp();\n float acc=bias1[lane];\n #pragma unroll\n for(int k=0;k<C_DIM;k++) acc=fmaf(w1[lane*C_DIM+k], s[k], acc);\n float* y1=h1+warp*H_DIM;\n y1[lane]=acc>0.0f?acc:0.0f;\n __syncwarp();\n float acc2=bias2[lane];\n #pragma unroll\n for(int k=0;k<H_DIM;k++) acc2=fmaf(w2[lane*H_DIM+k], y1[k], acc2);\n float* y2=h2+warp*H_DIM;\n y2[lane]=acc2>0.0f?acc2:0.0f;\n __syncwarp();\n float zsum=warp_sum(w3[lane]*y2[lane]);\n if(lane==0){\n float p=1.0f/(1.0f+expf(-(zsum+bias3[0])));\n long o1=((long)b*N+i)*N+j;\n long o2=((long)b*N+j)*N+i;\n out[o1]=p;\n out[o2]=p;\n }\n}\n\ntorch::Tensor contact_forward(torch::Tensor pair_repr, torch::Tensor ln_weight, torch::Tensor ln_bias,\n torch::Tensor w1, torch::Tensor b1, torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3){\n int B=(int)pair_repr.size(0), N=(int)pair_repr.size(1);\n auto out=torch::empty({B,N,N}, pair_repr.options());\n int tri=N*(N+1)/2;\n int total_tri=B*tri;\n int blocks=(total_tri+WARPS-1)/WARPS;\n contact_kernel_tri<<<blocks,BLOCK,0,at::cuda::getCurrentCUDAStream()>>>(\n pair_repr.data_ptr<float>(), ln_weight.data_ptr<float>(), ln_bias.data_ptr<float>(),\n w1.data_ptr<float>(), b1.data_ptr<float>(), w2.data_ptr<float>(), b2.data_ptr<float>(),\n w3.data_ptr<float>(), b3.data_ptr<float>(), out.data_ptr<float>(), N, tri, total_tri);\n return out;\n}\n\"\"\"\n_ext=load_inline(name=\"contact_head_ext_tri1\", cpp_sources=cpp_source, cuda_sources=cuda_source,\n functions=[\"contact_forward\"], extra_cuda_cflags=[\"-O3\",\"--use_fast_math\"], extra_cflags=[\"-O3\"], verbose=False)\nclass ModelNew(nn.Module):\n def __init__(self, pair_dim, hidden_dim):\n super().__init__()\n self.layer_norm=nn.LayerNorm(pair_dim)\n self.mlp=nn.Sequential(nn.Linear(pair_dim,hidden_dim),nn.ReLU(),nn.Linear(hidden_dim,hidden_dim),nn.ReLU(),nn.Linear(hidden_dim,1))\n def forward(self,pair_repr:torch.Tensor)->torch.Tensor:\n return _ext.contact_forward(pair_repr,self.layer_norm.weight,self.layer_norm.bias,\n self.mlp[0].weight,self.mlp[0].bias,self.mlp[2].weight,self.mlp[2].bias,self.mlp[4].weight,self.mlp[4].bias)\n"
}