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 = r'''\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n__global__ void importance_sampling_kernel(\n const float* samples,\n const float* mix_mus,\n const float* mix_stds,\n const float* log_mix_weights,\n const float* q_mu,\n const float* q_std,\n float* log_weights,\n int num_samples, int dim, int num_components\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (idx < num_samples) {\n float max_log_comp = -1e30f;\n float sum_exp = 0.0f;\n \n for (int k = 0; k < num_components; k++) {\n float log_comp = log_mix_weights[k];\n for (int d = 0; d < dim; d++) {\n float diff = samples[idx * dim + d] - mix_mus[k * dim + d];\n float std = mix_stds[k * dim + d];\n float z = diff / std;\n log_comp += -0.5f * z * z - logf(std);\n }\n \n if (log_comp > max_log_comp) {\n sum_exp = expf(max_log_comp - log_comp) * sum_exp + 1.0f;\n max_log_comp = log_comp;\n } else {\n sum_exp += expf(log_comp - max_log_comp);\n }\n }\n \n float log_p = max_log_comp + logf(sum_exp + 1e-30f);\n \n float log_q = 0.0f;\n for (int d = 0; d < dim; d++) {\n float diff = samples[idx * dim + d] - q_mu[d];\n float std = q_std[d];\n float z = diff / std;\n log_q += -0.5f * z * z - logf(std);\n }\n \n log_weights[idx] = log_p - log_q;\n }\n}\n\n__global__ void logsumexp_kernel(\n const float* log_weights,\n float* result,\n int num_samples\n) {\n __shared__ float max_val;\n \n int tid = threadIdx.x;\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n \n float local_max = -1e30f;\n \n for (int i = idx; i < num_samples; i += blockDim.x) {\n local_max = fmaxf(local_max, log_weights[i]);\n }\n \n if (tid == 0) max_val = -1e30f;\n __syncthreads();\n \n for (int offset = 16; offset > 0; offset >>= 1) {\n local_max = fmaxf(local_max, __shfl_down_sync(0xffffffff, local_max, offset));\n }\n \n if (tid == 0) max_val = local_max;\n __syncthreads();\n \n float global_max = max_val;\n \n float local_sum = 0.0f;\n for (int i = idx; i < num_samples; i += blockDim.x) {\n local_sum += expf(log_weights[i] - global_max);\n }\n \n for (int offset = 16; offset > 0; offset >>= 1) {\n local_sum += __shfl_down_sync(0xffffffff, local_sum, offset);\n }\n \n if (tid == 0) {\n result[0] = global_max + logf(local_sum + 1e-30f);\n }\n}\n\n__global__ void normalize_and_reduce_kernel(\n const float* log_weights,\n const float* samples,\n float* weighted_mean,\n float* ess_val,\n float log_sum_weights,\n int num_samples, int dim\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (idx < num_samples) {\n float w = expf(log_weights[idx] - log_sum_weights);\n float w_sq = w * w;\n \n for (int d = 0; d < dim; d++) {\n atomicAdd(&weighted_mean[d], w * samples[idx * dim + d]);\n }\n atomicAdd(ess_val, w_sq);\n }\n}\n\nvoid run_importance_sampling_kernel(\n const float* samples, const float* mix_mus, const float* mix_stds,\n const float* log_mix_weights, const float* q_mu, const float* q_std,\n float* log_weights, int num_samples, int dim, int num_components\n) {\n const int block_size = 256;\n const int num_blocks = (num_samples + block_size - 1) / block_size;\n importance_sampling_kernel<<<num_blocks, block_size>>>(\n samples, mix_mus, mix_stds, log_mix_weights, q_mu, q_std, log_weights,\n num_samples, dim, num_components\n );\n}\n\nvoid run_logsumexp_kernel(const float* log_weights, float* result, int num_samples) {\n const int block_size = 256;\n logsumexp_kernel<<<1, block_size>>>(log_weights, result, num_samples);\n}\n\nvoid run_normalize_and_reduce_kernel(\n const float* log_weights, const float* samples, float* weighted_mean,\n float* ess_val, float log_sum_weights, int num_samples, int dim\n) {\n const int block_size = 256;\n const int num_blocks = (num_samples + block_size - 1) / block_size;\n normalize_and_reduce_kernel<<<num_blocks, block_size>>>(\n log_weights, samples, weighted_mean, ess_val, log_sum_weights, num_samples, dim\n );\n}\n'''\n\ncpp_source = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nvoid run_importance_sampling_kernel(\n const float* samples, const float* mix_mus, const float* mix_stds,\n const float* log_mix_weights, const float* q_mu, const float* q_std,\n float* log_weights, int num_samples, int dim, int num_components);\n\nvoid run_logsumexp_kernel(const float* log_weights, float* result, int num_samples);\n\nvoid run_normalize_and_reduce_kernel(\n const float* log_weights, const float* samples, float* weighted_mean,\n float* ess_val, float log_sum_weights, int num_samples, int dim);\n\ntorch::Tensor importance_sampling_forward(\n torch::Tensor samples,\n torch::Tensor mix_mus,\n torch::Tensor mix_stds,\n torch::Tensor mix_weights,\n torch::Tensor q_mu,\n torch::Tensor q_std\n) {\n int num_samples = samples.size(0);\n int dim = samples.size(1);\n int num_components = mix_weights.size(0);\n \n auto log_weights = torch::empty({num_samples}, samples.options());\n auto log_mix_weights = torch::log(mix_weights);\n auto weighted_mean = torch::zeros({dim}, samples.options());\n auto ess_val = torch::zeros({1}, samples.options());\n auto log_sum_weights = torch::empty({1}, samples.options());\n \n run_importance_sampling_kernel(\n samples.data_ptr<float>(),\n mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(),\n log_mix_weights.data_ptr<float>(),\n q_mu.data_ptr<float>(),\n q_std.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n num_samples, dim, num_components\n );\n \n run_logsumexp_kernel(\n log_weights.data_ptr<float>(),\n log_sum_weights.data_ptr<float>(),\n num_samples\n );\n \n float log_sum;\n cudaMemcpy(&log_sum, log_sum_weights.data_ptr<float>(), sizeof(float), cudaMemcpyDeviceToHost);\n \n run_normalize_and_reduce_kernel(\n log_weights.data_ptr<float>(),\n samples.data_ptr<float>(),\n weighted_mean.data_ptr<float>(),\n ess_val.data_ptr<float>(),\n log_sum,\n num_samples, dim\n );\n \n float ess_sum;\n cudaMemcpy(&ess_sum, ess_val.data_ptr<float>(), sizeof(float), cudaMemcpyDeviceToHost);\n ess_sum = 1.0f / ess_sum;\n cudaMemcpy(ess_val.data_ptr<float>(), &ess_sum, sizeof(float), cudaMemcpyHostToDevice);\n \n return torch::cat({weighted_mean, ess_val});\n}\n'''\n\nmodule = load_inline(\n name=\"importance_sampling_ext\",\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=[\"importance_sampling_forward\"],\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 result = module.importance_sampling_forward(\n samples, self.mix_mus, self.mix_stds, self.mix_weights, self.q_mu, self.q_std\n )\n weighted_mean = result[:self.dim]\n ess = result[self.dim]\n return weighted_mean, ess"
}