call → compile_kernel
callcompile_kernel
{
"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>\n#include <vector>\nstd::vector<torch::Tensor> snis_forward_cuda(torch::Tensor samples, torch::Tensor mix_weights,\n torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor q_mu, torch::Tensor q_std);\nstd::vector<torch::Tensor> snis_forward(torch::Tensor samples, torch::Tensor mix_weights,\n torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor q_mu, torch::Tensor q_std) {\n return snis_forward_cuda(samples, mix_weights, mix_mus, mix_stds, q_mu, q_std);\n}\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n#include <cmath>\n\n__device__ __forceinline__ float compute_logw_16_3(const float* __restrict__ x,\n const float* __restrict__ mix_weights,\n const float* __restrict__ mu) {\n float sx2 = 0.0f;\n float d0 = 0.0f, d1 = 0.0f, d2 = 0.0f;\n #pragma unroll\n for (int d = 0; d < 16; ++d) {\n float xd = x[d];\n sx2 += xd * xd;\n float a0 = xd - mu[d];\n float a1 = xd - mu[16 + d];\n float a2 = xd - mu[32 + d];\n d0 += a0 * a0;\n d1 += a1 * a1;\n d2 += a2 * a2;\n }\n float lw = logf(mix_weights[0]);\n float c0 = -0.5f * d0 + lw;\n float c1 = -0.5f * d1 + lw;\n float c2 = -0.5f * d2 + lw;\n float m = fmaxf(c0, fmaxf(c1, c2));\n float lse = m + logf(expf(c0 - m) + expf(c1 - m) + expf(c2 - m));\n return lse + 0.05555555555555556f * sx2 + 17.577796618689757f;\n}\n\n__global__ void logw16_kernel(const float* __restrict__ samples,\n const float* __restrict__ mix_weights,\n const float* __restrict__ mix_mus,\n int N,\n float* __restrict__ logw,\n float* __restrict__ partial_max) {\n __shared__ float s[256];\n int tid = threadIdx.x;\n int i = blockIdx.x * 256 + tid;\n float v = -INFINITY;\n if (i < N) {\n v = compute_logw_16_3(samples + ((long long)i << 4), mix_weights, mix_mus);\n logw[i] = v;\n }\n s[tid] = v;\n __syncthreads();\n #pragma unroll\n for (int stride = 128; stride > 0; stride >>= 1) {\n if (tid < stride) s[tid] = fmaxf(s[tid], s[tid + stride]);\n __syncthreads();\n }\n if (tid == 0) partial_max[blockIdx.x] = s[0];\n}\n\n__global__ void reduce16_kernel(const float* __restrict__ samples,\n const float* __restrict__ logw,\n const float* __restrict__ partial_max,\n int N, int num_blocks,\n float* __restrict__ out_mean,\n float* __restrict__ out_ess) {\n extern __shared__ float smem[];\n int tid = threadIdx.x;\n float* s_max = smem;\n float* s_sum = smem + 256;\n float* s_sum2 = smem + 512;\n float* s_num = smem + 768;\n\n float lm = -INFINITY;\n for (int j = tid; j < num_blocks; j += 256) lm = fmaxf(lm, partial_max[j]);\n s_max[tid] = lm;\n __syncthreads();\n #pragma unroll\n for (int stride = 128; stride > 0; stride >>= 1) {\n if (tid < stride) s_max[tid] = fmaxf(s_max[tid], s_max[tid + stride]);\n __syncthreads();\n }\n float maxlw = s_max[0];\n\n float lsum = 0.0f, lsum2 = 0.0f;\n float n0=0.0f,n1=0.0f,n2=0.0f,n3=0.0f,n4=0.0f,n5=0.0f,n6=0.0f,n7=0.0f;\n float n8=0.0f,n9=0.0f,n10=0.0f,n11=0.0f,n12=0.0f,n13=0.0f,n14=0.0f,n15=0.0f;\n for (int i = tid; i < N; i += 256) {\n float a = expf(logw[i] - maxlw);\n lsum += a;\n lsum2 += a * a;\n const float* x = samples + ((long long)i << 4);\n n0 += a*x[0]; n1 += a*x[1]; n2 += a*x[2]; n3 += a*x[3];\n n4 += a*x[4]; n5 += a*x[5]; n6 += a*x[6]; n7 += a*x[7];\n n8 += a*x[8]; n9 += a*x[9]; n10 += a*x[10]; n11 += a*x[11];\n n12 += a*x[12]; n13 += a*x[13]; n14 += a*x[14]; n15 += a*x[15];\n }\n s_sum[tid] = lsum;\n s_sum2[tid] = lsum2;\n s_num[0*256+tid]=n0; s_num[1*256+tid]=n1; s_num[2*256+tid]=n2; s_num[3*256+tid]=n3;\n s_num[4*256+tid]=n4; s_num[5*256+tid]=n5; s_num[6*256+tid]=n6; s_num[7*256+tid]=n7;\n s_num[8*256+tid]=n8; s_num[9*256+tid]=n9; s_num[10*256+tid]=n10; s_num[11*256+tid]=n11;\n s_num[12*256+tid]=n12; s_num[13*256+tid]=n13; s_num[14*256+tid]=n14; s_num[15*256+tid]=n15;\n __syncthreads();\n\n #pragma unroll\n for (int stride = 128; stride > 0; stride >>= 1) {\n if (tid < stride) {\n s_sum[tid] += s_sum[tid + stride];\n s_sum2[tid] += s_sum2[tid + stride];\n #pragma unroll\n for (int d = 0; d < 16; ++d) s_num[d*256 + tid] += s_num[d*256 + tid + stride];\n }\n __syncthreads();\n }\n if (tid == 0) {\n float inv = 1.0f / s_sum[0];\n #pragma unroll\n for (int d = 0; d < 16; ++d) out_mean[d] = s_num[d*256] * inv;\n out_ess[0] = s_sum[0] * s_sum[0] / s_sum2[0];\n }\n}\n\nstd::vector<torch::Tensor> snis_forward_cuda(torch::Tensor samples, torch::Tensor mix_weights,\n torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor q_mu, torch::Tensor q_std) {\n samples = samples.contiguous();\n int N = (int)samples.size(0);\n auto mean = torch::empty({16}, samples.options());\n auto ess = torch::empty({}, samples.options());\n auto logw = torch::empty({N}, samples.options());\n const int threads = 256;\n int blocks = (N + threads - 1) / threads;\n auto partial = torch::empty({blocks}, samples.options());\n logw16_kernel<<<blocks, threads>>>(samples.data_ptr<float>(), mix_weights.data_ptr<float>(), mix_mus.data_ptr<float>(),\n N, logw.data_ptr<float>(), partial.data_ptr<float>());\n reduce16_kernel<<<1, threads, 19 * threads * sizeof(float)>>>(samples.data_ptr<float>(), logw.data_ptr<float>(),\n partial.data_ptr<float>(), N, blocks,\n mean.data_ptr<float>(), ess.data_ptr<float>());\n return {mean, ess};\n}\n\"\"\"\n\n_ext = load_inline(name=\"snis_ext_v3\", cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=[\"snis_forward\"], verbose=False, extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"])\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_components=3):\n super().__init__()\n self.dim = dim\n self.num_components = num_components\n self.register_buffer(\"mix_weights\", torch.ones(num_components) / num_components)\n mus = torch.randn(num_components, dim) * 2.0\n self.register_buffer(\"mix_mus\", mus)\n self.register_buffer(\"mix_stds\", torch.ones(num_components, dim))\n self.register_buffer(\"q_mu\", torch.zeros(dim))\n self.register_buffer(\"q_std\", torch.ones(dim) * 3.0)\n def forward(self, samples: torch.Tensor) -> tuple:\n out = _ext.snis_forward(samples, self.mix_weights, self.mix_mus, self.mix_stds, self.q_mu, self.q_std)\n return (out[0], out[1])\n"
}