call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r\"\"\"\n#include <torch/extension.h>\ntorch::Tensor gibbs_forward(torch::Tensor x, torch::Tensor precision);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/ATen.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__device__ __forceinline__ const float* selptr(int i,\n const float* n0,const float* n1,const float* n2,const float* n3,const float* n4,const float* n5,const float* n6,const float* n7,\n const float* n8,const float* n9,const float* n10,const float* n11,const float* n12,const float* n13,const float* n14,const float* n15,\n const float* n16,const float* n17,const float* n18,const float* n19,const float* n20,const float* n21,const float* n22,const float* n23,\n const float* n24,const float* n25,const float* n26,const float* n27,const float* n28,const float* n29,const float* n30,const float* n31) {\n switch(i){\n case 0:return n0; case 1:return n1; case 2:return n2; case 3:return n3; case 4:return n4; case 5:return n5; case 6:return n6; case 7:return n7;\n case 8:return n8; case 9:return n9; case 10:return n10; case 11:return n11; case 12:return n12; case 13:return n13; case 14:return n14; case 15:return n15;\n case 16:return n16; case 17:return n17; case 18:return n18; case 19:return n19; case 20:return n20; case 21:return n21; case 22:return n22; case 23:return n23;\n case 24:return n24; case 25:return n25; case 26:return n26; case 27:return n27; case 28:return n28; case 29:return n29; case 30:return n30; default:return n31;\n }\n}\n\n__global__ __launch_bounds__(128, 2)\nvoid gibbs32_kernel(const float* __restrict__ x, float* __restrict__ y, const float* __restrict__ p, int B,\n const float* n0,const float* n1,const float* n2,const float* n3,const float* n4,const float* n5,const float* n6,const float* n7,\n const float* n8,const float* n9,const float* n10,const float* n11,const float* n12,const float* n13,const float* n14,const float* n15,\n const float* n16,const float* n17,const float* n18,const float* n19,const float* n20,const float* n21,const float* n22,const float* n23,\n const float* n24,const float* n25,const float* n26,const float* n27,const float* n28,const float* n29,const float* n30,const float* n31) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n const float* in = x + ((long long)b)*32;\n float v0=in[0],v1=in[1],v2=in[2],v3=in[3],v4=in[4],v5=in[5],v6=in[6],v7=in[7];\n float v8=in[8],v9=in[9],v10=in[10],v11=in[11],v12=in[12],v13=in[13],v14=in[14],v15=in[15];\n float v16=in[16],v17=in[17],v18=in[18],v19=in[19],v20=in[20],v21=in[21],v22=in[22],v23=in[23];\n float v24=in[24],v25=in[25],v26=in[26],v27=in[27],v28=in[28],v29=in[29],v30=in[30],v31=in[31];\n float vals[32];\n vals[0]=v0; vals[1]=v1; vals[2]=v2; vals[3]=v3; vals[4]=v4; vals[5]=v5; vals[6]=v6; vals[7]=v7;\n vals[8]=v8; vals[9]=v9; vals[10]=v10; vals[11]=v11; vals[12]=v12; vals[13]=v13; vals[14]=v14; vals[15]=v15;\n vals[16]=v16; vals[17]=v17; vals[18]=v18; vals[19]=v19; vals[20]=v20; vals[21]=v21; vals[22]=v22; vals[23]=v23;\n vals[24]=v24; vals[25]=v25; vals[26]=v26; vals[27]=v27; vals[28]=v28; vals[29]=v29; vals[30]=v30; vals[31]=v31;\n #pragma unroll 32\n for (int i=0;i<32;i++) {\n const float* row = p + i*32;\n float sum = 0.0f;\n #pragma unroll 32\n for (int j=0;j<32;j++) {\n if (j != i) sum += vals[j] * row[j];\n }\n float pii = row[i];\n const float* np = selptr(i,n0,n1,n2,n3,n4,n5,n6,n7,n8,n9,n10,n11,n12,n13,n14,n15,n16,n17,n18,n19,n20,n21,n22,n23,n24,n25,n26,n27,n28,n29,n30,n31);\n vals[i] = -sum / pii + rsqrtf(pii) * np[b];\n }\n float* out = y + ((long long)b)*32;\n #pragma unroll\n for(int j=0;j<32;j++) out[j]=vals[j];\n}\n\ntorch::Tensor gibbs_forward(torch::Tensor x, torch::Tensor precision) {\n auto y = torch::empty_like(x);\n int B = (int)x.size(0);\n auto n0=torch::randn({B}, x.options()); auto n1=torch::randn({B}, x.options()); auto n2=torch::randn({B}, x.options()); auto n3=torch::randn({B}, x.options());\n auto n4=torch::randn({B}, x.options()); auto n5=torch::randn({B}, x.options()); auto n6=torch::randn({B}, x.options()); auto n7=torch::randn({B}, x.options());\n auto n8=torch::randn({B}, x.options()); auto n9=torch::randn({B}, x.options()); auto n10=torch::randn({B}, x.options()); auto n11=torch::randn({B}, x.options());\n auto n12=torch::randn({B}, x.options()); auto n13=torch::randn({B}, x.options()); auto n14=torch::randn({B}, x.options()); auto n15=torch::randn({B}, x.options());\n auto n16=torch::randn({B}, x.options()); auto n17=torch::randn({B}, x.options()); auto n18=torch::randn({B}, x.options()); auto n19=torch::randn({B}, x.options());\n auto n20=torch::randn({B}, x.options()); auto n21=torch::randn({B}, x.options()); auto n22=torch::randn({B}, x.options()); auto n23=torch::randn({B}, x.options());\n auto n24=torch::randn({B}, x.options()); auto n25=torch::randn({B}, x.options()); auto n26=torch::randn({B}, x.options()); auto n27=torch::randn({B}, x.options());\n auto n28=torch::randn({B}, x.options()); auto n29=torch::randn({B}, x.options()); auto n30=torch::randn({B}, x.options()); auto n31=torch::randn({B}, x.options());\n int threads=128; int blocks=(B+threads-1)/threads;\n gibbs32_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(x.data_ptr<float>(), y.data_ptr<float>(), precision.data_ptr<float>(), B,\n n0.data_ptr<float>(),n1.data_ptr<float>(),n2.data_ptr<float>(),n3.data_ptr<float>(),n4.data_ptr<float>(),n5.data_ptr<float>(),n6.data_ptr<float>(),n7.data_ptr<float>(),\n n8.data_ptr<float>(),n9.data_ptr<float>(),n10.data_ptr<float>(),n11.data_ptr<float>(),n12.data_ptr<float>(),n13.data_ptr<float>(),n14.data_ptr<float>(),n15.data_ptr<float>(),\n n16.data_ptr<float>(),n17.data_ptr<float>(),n18.data_ptr<float>(),n19.data_ptr<float>(),n20.data_ptr<float>(),n21.data_ptr<float>(),n22.data_ptr<float>(),n23.data_ptr<float>(),\n n24.data_ptr<float>(),n25.data_ptr<float>(),n26.data_ptr<float>(),n27.data_ptr<float>(),n28.data_ptr<float>(),n29.data_ptr<float>(),n30.data_ptr<float>(),n31.data_ptr<float>());\n return y;\n}\n\"\"\"\n_ext = load_inline(name=\"gibbs_ext_v5\", cpp_sources=_cpp_src, cuda_sources=_cuda_src, functions=[\"gibbs_forward\"], extra_cuda_cflags=[\"-O3\"], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim):\n super().__init__(); self.dim=dim\n L = torch.randn(dim, dim) * 0.1\n precision = L @ L.t() + torch.eye(dim) * 2.0\n self.register_buffer(\"precision\", precision)\n self.register_buffer(\"mu\", torch.zeros(dim))\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.gibbs_forward(x, self.precision)\n"
}