turn 132.7s turn · 32.7s LLM · 0.0s tools · tokens in=2,542 · out=1,962 · total=4,582
reasoning
H100, high bandwidth, so memory bound probably.
assistant```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
__device__ float logsumexp_device(float* vals, int n) {
float maxv = -INFINITY;
for (int i = 0; i < n; i++) {
if (vals[i] > maxv) maxv = vals[i];
}
float sum = 0.0f;
for (int i = 0; i < n; i++) {
sum += expf(vals[i] - maxv);
}
return maxv + logf(sum);
}
__global__ void importance_sampling_kernel(
const float* __restrict__ samples,
const float* __restrict__ mix_mus,
const float* __restrict__ mix_stds,
const float* __restrict__ mix_weights_log,
const float* __restrict__ q_mu,
const float* __restrict__ q_std_log,
float* __restrict__ log_w_out,
float* __restrict__ weighted_sum,
float* __restrict__ sum_w2,
int N,
int D,
int K
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
// Compute log_p
float log_comp[8]; // assume K<=8
for (int k = 0; k < K; k++) {
float sum = 0.0f;
for (int d = 0; d < D; d++) {
float diff = samples[i * D + d] - mix_mus[k * D + d];
float std = mix_stds[k * D + d];
sum += -0.5f * (diff / std) * (diff / std) - logf(std);
}
log_comp[k] = sum + mix_weights_log[k];
}
float log_p = logsumexp_device(log_comp, K);
// Compute log_q
float sum_q = 0.0f;
for (int d = 0; d < D; d++) {
float diff = samples[i * D + d] - q_mu[d];
float std_log = q_std_log[d];
sum_q += -0.5f * (diff * diff) * expf(-2.0f * std_log) - std_log;
}
float log_q = sum_q;
float lw = log_p - log_q;
log_w_out[i] = lw;
// For reduction later, but we'll do in python for simplicity or separate
}
__global__ void reduce_kernel(
const float* __restrict__ log_w,
const float* __restrict__ samples,
float* __restrict__ weighted_mean,
float* __restrict__ ess_out,
int N,
int D
) {
extern __shared__ float sdata[];
float max_log = -INFINITY;
for (int i = threadIdx.x; i < N; i += blockDim.x) {
if (log_w[i] > max_log) max_log = log_w[i];
}
// block reduce max
sdata[threadIdx.x] = max_log;
__syncthreads();
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (threadIdx.x < s) {
sdata[threadIdx.x] = fmaxf(sdata[threadIdx.x], sdata[threadIdx.x + s]);
}
__syncthreads();
}
max_log = sdata[0];
float sum_exp = 0.0f;
float sum_w2 = 0.0f;
for (int i = threadIdx.x; i < N; i += blockDim.x) {
float w = expf(log_w[i] - max_log);
sum_exp += w;
sum_w2 += w * w;
}
sdata[threadIdx.x] = sum_exp;
__syncthreads();
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (threadIdx.x < s) sdata[threadIdx.x] += sdata[threadIdx.x + s];
__syncthreads();
}
sum_exp = sdata[0];
sdata[threadIdx.x] = sum_w2;
__syncthreads();
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (threadIdx.x < s) sdata[threadIdx.x] += sdata[threadIdx.x + s];
__syncthreads();
}
sum_w2 = sdata[0];
float log_sum = max_log + logf(sum_exp);
float inv_sum = 1.0f / sum_exp;
if (threadIdx.x < D) {
float wsum = 0.0f;
for (int i = 0; i < N; i++) {
float w = expf(log_w[i] - log_sum);
wsum += w * samples[i * D + threadIdx.x];
}
weighted_mean[threadIdx.x] = wsum;
}
if (threadIdx.x == 0) {
ess_out[0] = 1.0f / sum_w2;
}
}
torch::Tensor compute_log_w(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds, torch::Tensor mix_weights, torch::Tensor q_mu, torch::Tensor q_std) {
int N = samples.size(0);
int D = samples.size(1);
int K = mix_mus.size(0);
auto log_w = torch::empty({N}, samples.options());
float mix_weights_log[8];
auto mw = mix_weights.cpu();
for (int k=0; k<K; k++) mix_weights_log[k] = logf(mw[k].item<float>());
auto mix_weights_log_t = torch::from_blob(mix_weights_log, {K}, torch::kFloat32).to(samples.device());
float q_std_log[32];
auto qs = q_std.cpu();
for (int d=0; d<D; d++) q_std_log[d] = logf(qs[d].item<float>());
auto q_std_log_t = torch::from_blob(q_std_log, {D}, torch::kFloat32).to(samples.device());
dim3 block(256);
dim3 grid((N + 255) / 256);
importance_sampling_kernel<<<grid, block>>>(
samples.data_ptr<float>(),
mix_mus.data_ptr<float>(),
mix_stds.data_ptr<float>(),
mix_weights_log_t.data_ptr<float>(),
q_mu.data_ptr<float>(),
q_std_log_t.data_ptr<float>(),
log_w.data_ptr<float>(),
nullptr, nullptr,
N, D, K
);
return log_w;
}
std::vector<torch::Tensor> finalize(torch::Tensor log_w, torch::Tensor samples) {
int N = samples.size(0);
int D = samples.size(1);
auto weighted_mean = torch::empty({D}, samples.options());
auto ess = torch::empty({1}, samples.options());
dim3 block(256);
reduce_kernel<<<1, block, 256 * sizeof(float)>>>(
log_w.data_ptr<float>(),
samples.data_ptr<float>(),
weighted_mean.data_ptr<float>(),
ess.data_ptr<float>(),
N, D
);
return {weighted_mean, ess.squeeze()};
}
'''
cpp_source = '''
torch::Tensor compute_log_w(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds, torch::Tensor mix_weights, torch::Tensor q_mu, torch::Tensor q_std);
std::vector<torch::Tensor> finalize(torch::Tensor log_w, torch::Tensor samples);
'''
module = load_inline(
name='importance_sampling',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['compute_log_w', 'finalize'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, dim, num_components=3):
super().__init__()
self.dim = dim
self.num_components = num_components
self.register_buffer("mix_weights", torch.ones(num_components) / num_components)
mus = torch.randn(num_components, dim) * 2.0
self.register_buffer("mix_mus", mus)
self.register_buffer("mix_stds", torch.ones(num_components, dim))
self.register_buffer("q_mu", torch.zeros(dim))
self.register_buffer("q_std", torch.ones(dim) * 3.0)
def forward(self, samples: torch.Tensor) -> tuple:
log_w = module.compute_log_w(samples, self.mix_mus, self.mix_stds, self.mix_weights, self.q_mu, self.q_std)
return module.finalize(log_w, samples)
```