call → submit_kernel
callsubmit_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>\n\nstd::vector<torch::Tensor> snis_forward_cuda(torch::Tensor samples,\n torch::Tensor mix_weights,\n torch::Tensor mix_mus,\n torch::Tensor mix_stds,\n torch::Tensor q_mu,\n torch::Tensor q_std);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <vector>\n#include <cfloat>\n#include <cmath>\n\n__global__ void logw_max_d16k3_kernel(const float* __restrict__ x,\n const float* __restrict__ mix_w,\n const float* __restrict__ mus,\n const float* __restrict__ stds,\n const float* __restrict__ q_mu,\n const float* __restrict__ q_std,\n float* __restrict__ logw,\n float* __restrict__ max_out,\n int N) {\n __shared__ float smem[256];\n int tid = threadIdx.x;\n float local_max = -FLT_MAX;\n\n for (int i = tid; i < N; i += blockDim.x) {\n const float* xi = x + ((long long)i) * 16;\n float xv0 = xi[0], xv1 = xi[1], xv2 = xi[2], xv3 = xi[3];\n float xv4 = xi[4], xv5 = xi[5], xv6 = xi[6], xv7 = xi[7];\n float xv8 = xi[8], xv9 = xi[9], xv10 = xi[10], xv11 = xi[11];\n float xv12 = xi[12], xv13 = xi[13], xv14 = xi[14], xv15 = xi[15];\n\n float qlog = 0.0f;\n float d;\n d=(xv0-q_mu[0])/q_std[0]; qlog += -0.5f*d*d - logf(q_std[0]);\n d=(xv1-q_mu[1])/q_std[1]; qlog += -0.5f*d*d - logf(q_std[1]);\n d=(xv2-q_mu[2])/q_std[2]; qlog += -0.5f*d*d - logf(q_std[2]);\n d=(xv3-q_mu[3])/q_std[3]; qlog += -0.5f*d*d - logf(q_std[3]);\n d=(xv4-q_mu[4])/q_std[4]; qlog += -0.5f*d*d - logf(q_std[4]);\n d=(xv5-q_mu[5])/q_std[5]; qlog += -0.5f*d*d - logf(q_std[5]);\n d=(xv6-q_mu[6])/q_std[6]; qlog += -0.5f*d*d - logf(q_std[6]);\n d=(xv7-q_mu[7])/q_std[7]; qlog += -0.5f*d*d - logf(q_std[7]);\n d=(xv8-q_mu[8])/q_std[8]; qlog += -0.5f*d*d - logf(q_std[8]);\n d=(xv9-q_mu[9])/q_std[9]; qlog += -0.5f*d*d - logf(q_std[9]);\n d=(xv10-q_mu[10])/q_std[10]; qlog += -0.5f*d*d - logf(q_std[10]);\n d=(xv11-q_mu[11])/q_std[11]; qlog += -0.5f*d*d - logf(q_std[11]);\n d=(xv12-q_mu[12])/q_std[12]; qlog += -0.5f*d*d - logf(q_std[12]);\n d=(xv13-q_mu[13])/q_std[13]; qlog += -0.5f*d*d - logf(q_std[13]);\n d=(xv14-q_mu[14])/q_std[14]; qlog += -0.5f*d*d - logf(q_std[14]);\n d=(xv15-q_mu[15])/q_std[15]; qlog += -0.5f*d*d - logf(q_std[15]);\n\n float c0 = logf(mix_w[0]);\n float c1 = logf(mix_w[1]);\n float c2 = logf(mix_w[2]);\n #pragma unroll\n for (int j = 0; j < 16; ++j) {\n float xx;\n if (j==0) xx=xv0; else if (j==1) xx=xv1; else if (j==2) xx=xv2; else if (j==3) xx=xv3;\n else if (j==4) xx=xv4; else if (j==5) xx=xv5; else if (j==6) xx=xv6; else if (j==7) xx=xv7;\n else if (j==8) xx=xv8; else if (j==9) xx=xv9; else if (j==10) xx=xv10; else if (j==11) xx=xv11;\n else if (j==12) xx=xv12; else if (j==13) xx=xv13; else if (j==14) xx=xv14; else xx=xv15;\n float s0 = stds[j];\n float t = (xx - mus[j]) / s0; c0 += -0.5f*t*t - logf(s0);\n float s1 = stds[16 + j];\n t = (xx - mus[16 + j]) / s1; c1 += -0.5f*t*t - logf(s1);\n float s2 = stds[32 + j];\n t = (xx - mus[32 + j]) / s2; c2 += -0.5f*t*t - logf(s2);\n }\n float m = fmaxf(c0, fmaxf(c1, c2));\n float lp = m + logf(expf(c0 - m) + expf(c1 - m) + expf(c2 - m));\n float lw = lp - qlog;\n logw[i] = lw;\n local_max = fmaxf(local_max, lw);\n }\n smem[tid] = local_max;\n __syncthreads();\n for (int s = blockDim.x >> 1; s > 0; s >>= 1) {\n if (tid < s) smem[tid] = fmaxf(smem[tid], smem[tid + s]);\n __syncthreads();\n }\n if (tid == 0) max_out[0] = smem[0];\n}\n\n__global__ void reduce_weighted_d16_kernel(const float* __restrict__ x,\n const float* __restrict__ logw,\n const float* __restrict__ max_in,\n float* __restrict__ mean,\n float* __restrict__ ess,\n int N) {\n extern __shared__ float sh[];\n int tid = threadIdx.x;\n const int BD = 256;\n float m = max_in[0];\n float sw = 0.0f, sw2 = 0.0f;\n float sx[16];\n #pragma unroll\n for (int j = 0; j < 16; ++j) sx[j] = 0.0f;\n\n for (int i = tid; i < N; i += blockDim.x) {\n float w = expf(logw[i] - m);\n sw += w;\n sw2 += w * w;\n const float* xi = x + ((long long)i) * 16;\n #pragma unroll\n for (int j = 0; j < 16; ++j) sx[j] += w * xi[j];\n }\n sh[tid] = sw;\n sh[BD + tid] = sw2;\n #pragma unroll\n for (int j = 0; j < 16; ++j) sh[(2 + j) * BD + tid] = sx[j];\n __syncthreads();\n\n for (int s = blockDim.x >> 1; s > 0; s >>= 1) {\n if (tid < s) {\n sh[tid] += sh[tid + s];\n sh[BD + tid] += sh[BD + tid + s];\n #pragma unroll\n for (int j = 0; j < 16; ++j) sh[(2 + j) * BD + tid] += sh[(2 + j) * BD + tid + s];\n }\n __syncthreads();\n }\n float inv = 1.0f / sh[0];\n if (tid < 16) mean[tid] = sh[(2 + tid) * BD] * inv;\n if (tid == 0) ess[0] = (sh[0] * sh[0]) / sh[BD];\n}\n\n__global__ void logw_max_generic_kernel(const float* __restrict__ x,\n const float* __restrict__ mix_w,\n const float* __restrict__ mus,\n const float* __restrict__ stds,\n const float* __restrict__ q_mu,\n const float* __restrict__ q_std,\n float* __restrict__ logw,\n float* __restrict__ max_out,\n int N, int D, int K) {\n __shared__ float smem[256];\n int tid = threadIdx.x;\n float local_max = -FLT_MAX;\n for (int i = tid; i < N; i += blockDim.x) {\n const float* xi = x + ((long long)i) * D;\n float qlog = 0.0f;\n for (int j = 0; j < D; ++j) {\n float t = (xi[j] - q_mu[j]) / q_std[j];\n qlog += -0.5f * t * t - logf(q_std[j]);\n }\n float maxc = -FLT_MAX;\n float sumexp = 0.0f;\n for (int k = 0; k < K; ++k) {\n float c = logf(mix_w[k]);\n for (int j = 0; j < D; ++j) {\n float st = stds[k * D + j];\n float t = (xi[j] - mus[k * D + j]) / st;\n c += -0.5f * t * t - logf(st);\n }\n if (k == 0) { maxc = c; sumexp = 1.0f; }\n else {\n if (c > maxc) { sumexp = sumexp * expf(maxc - c) + 1.0f; maxc = c; }\n else sumexp += expf(c - maxc);\n }\n }\n float lw = maxc + logf(sumexp) - qlog;\n logw[i] = lw;\n local_max = fmaxf(local_max, lw);\n }\n smem[tid] = local_max;\n __syncthreads();\n for (int s = blockDim.x >> 1; s > 0; s >>= 1) {\n if (tid < s) smem[tid] = fmaxf(smem[tid], smem[tid + s]);\n __syncthreads();\n }\n if (tid == 0) max_out[0] = smem[0];\n}\n\n__global__ void reduce_generic_kernel(const float* __restrict__ x,\n const float* __restrict__ logw,\n const float* __restrict__ max_in,\n float* __restrict__ mean,\n float* __restrict__ ess,\n int N, int D) {\n __shared__ float red[256];\n int tid = threadIdx.x;\n float m = max_in[0];\n // First reduce scalar sums.\n float sw = 0.0f, sw2 = 0.0f;\n for (int i = tid; i < N; i += blockDim.x) {\n float w = expf(logw[i] - m);\n sw += w; sw2 += w*w;\n }\n red[tid] = sw; __syncthreads();\n for (int s=blockDim.x>>1; s>0; s>>=1) { if (tid<s) red[tid]+=red[tid+s]; __syncthreads(); }\n float total = red[0];\n red[tid] = sw2; __syncthreads();\n for (int s=blockDim.x>>1; s>0; s>>=1) { if (tid<s) red[tid]+=red[tid+s]; __syncthreads(); }\n if (tid == 0) ess[0] = total*total/red[0];\n for (int d = 0; d < D; ++d) {\n float sx = 0.0f;\n for (int i = tid; i < N; i += blockDim.x) sx += expf(logw[i] - m) * x[(long long)i * D + d];\n red[tid] = sx; __syncthreads();\n for (int s=blockDim.x>>1; s>0; s>>=1) { if (tid<s) red[tid]+=red[tid+s]; __syncthreads(); }\n if (tid == 0) mean[d] = red[0] / total;\n __syncthreads();\n }\n}\n\nstd::vector<torch::Tensor> snis_forward_cuda(torch::Tensor samples,\n torch::Tensor mix_weights,\n torch::Tensor mix_mus,\n torch::Tensor mix_stds,\n torch::Tensor q_mu,\n torch::Tensor q_std) {\n TORCH_CHECK(samples.is_cuda(), \"samples must be CUDA\");\n TORCH_CHECK(samples.scalar_type() == torch::kFloat32, \"only float32 supported\");\n TORCH_CHECK(samples.is_contiguous(), \"samples must be contiguous\");\n int N = (int)samples.size(0);\n int D = (int)samples.size(1);\n int K = (int)mix_weights.size(0);\n auto opts = samples.options();\n auto logw = torch::empty({N}, opts);\n auto maxv = torch::empty({1}, opts);\n auto mean = torch::empty({D}, opts);\n auto ess = torch::empty({}, opts);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n const int threads = 256;\n if (D == 16 && K == 3) {\n logw_max_d16k3_kernel<<<1, threads, 0, stream>>>(\n samples.data_ptr<float>(), mix_weights.data_ptr<float>(), mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(), q_mu.data_ptr<float>(), q_std.data_ptr<float>(),\n logw.data_ptr<float>(), maxv.data_ptr<float>(), N);\n reduce_weighted_d16_kernel<<<1, threads, 18 * threads * sizeof(float), stream>>>(\n samples.data_ptr<float>(), logw.data_ptr<float>(), maxv.data_ptr<float>(),\n mean.data_ptr<float>(), ess.data_ptr<float>(), N);\n } else {\n logw_max_generic_kernel<<<1, threads, 0, stream>>>(\n samples.data_ptr<float>(), mix_weights.data_ptr<float>(), mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(), q_mu.data_ptr<float>(), q_std.data_ptr<float>(),\n logw.data_ptr<float>(), maxv.data_ptr<float>(), N, D, K);\n reduce_generic_kernel<<<1, threads, 0, stream>>>(\n samples.data_ptr<float>(), logw.data_ptr<float>(), maxv.data_ptr<float>(),\n mean.data_ptr<float>(), ess.data_ptr<float>(), N, D);\n }\n return {mean, ess};\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"snis_importance_ext_v1\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"snis_forward_cuda\"],\n extra_cflags=[\"-O3\"],\n extra_cuda_cflags=[\"-O3\"],\n verbose=False,\n)\n\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\n def forward(self, samples: torch.Tensor) -> tuple:\n out = _ext.snis_forward_cuda(samples.contiguous(), self.mix_weights.contiguous(),\n self.mix_mus.contiguous(), self.mix_stds.contiguous(),\n self.q_mu.contiguous(), self.q_std.contiguous())\n return (out[0], out[1])\n"
}