correct
16_ConvTranspose2d_Mish_Add_Hardtanh_Scaling.py
started
2026-05-05T00:28:52Z
timingkernelbench.timing
kernel mean / median / std
7.09 / — / 0.69 μs
ref mean / median / std
11.40 / — / 4.62 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
2.29x
energyNVML (pynvml)
ratio (ref / kernel)
0.67x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
13188.49 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Model that performs a transposed convolution, applies Mish activation, adds a value,
applies Hardtanh activation, and scales the output.
"""
def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, add_value, scale):
super(Model, self).__init__()
self.conv_transpose = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride, padding, output_padding)
self.add_value = add_value
self.scale = scale
def forward(self, x):
x = self.conv_transpose(x)
x = torch.nn.functional.mish(x) # Mish activation
x = x + self.add_value
x = torch.nn.functional.hardtanh(x, min_val=-1, max_val=1) # Hardtanh activation
x = x * self.scale # Scaling
return x
batch_size = 128
in_channels = 64
out_channels = 64
height = width = 128
kernel_size = 3
stride = 2
padding = 1
output_padding = 1
add_value = 0.5
scale = 2
def get_inputs():
return [torch.rand(batch_size, in_channels, height, width)]
def get_init_inputs():
return [in_channels, out_channels, kernel_size, stride, padding, output_padding, add_value, scale]turn 040.7s turn · 40.7s LLM · 0.0s tools · FINAL
reasoning
**Considering computation methods**
I’m thinking about whether I could use float4 if everything is aligned and the number of elements is a multiple. If that's the case, I might run into needing to apply computationally intensive functions like exp or tanh. This could get a bit heavy on the computation side. It’s an interesting trade-off to weigh the benefits of using float4 against the potential computational costs. I want to approach this thoughtfully!
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nfrom torch.utils.cpp_extension import load_inline\n\n\n_cpp_src = r\"\"\"\n#include <torch/extension.h>\n\ntorch::Tensor fused_mish_add_hardtanh_scale(torch::Tensor x, double add_value, double scale);\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\n__device__ __forceinline__ float softplus_stable(float v) {\n // Matches the usual stable softplus formulation used by PyTorch closely\n // enough for fp32 tolerance, while avoiding overflow for large positives.\n if (v > 20.0f) return v;\n if (v < -20.0f) return __expf(v);\n return log1pf(__expf(v));\n}\n\n__device__ __forceinline__ float op_one(float v, float addv, float scale) {\n // mish(v) = v * tanh(softplus(v))\n float sp = softplus_stable(v);\n float y = v * tanhf(sp);\n y += addv;\n // hardtanh with min=-1, max=1, then scale\n y = fminf(1.0f, fmaxf(-1.0f, y));\n return y * scale;\n}\n\n__global__ void fused_mish_add_hardtanh_scale_kernel(float* __restrict__ data,\n int64_t n,\n float addv,\n float scale) {\n int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;\n int64_t stride = (int64_t)blockDim.x * gridDim.x;\n for (; idx < n; idx += stride) {\n data[idx] = op_one(data[idx], addv, scale);\n }\n}\n\ntorch::Tensor fused_mish_add_hardtanh_scale(torch::Tensor x, double add_value, double scale) {\n const int64_t n = x.numel();\n if (n == 0) return x;\n const int threads = 256;\n int blocks = (int)((n + threads - 1) / threads);\n // Keep the grid reasonably sized; the grid-stride loop covers the rest.\n if (blocks > 65535) blocks = 65535;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n fused_mish_add_hardtanh_scale_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), n, (float)add_value, (float)scale);\n return x;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"fused_mish_add_hardtanh_scale_ext\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"fused_mish_add_hardtanh_scale\"],\n extra_cflags=[\"-O3\"],\n extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"],\n verbose=False,\n)\n\n\nclass ModelNew(nn.Module):\n \"\"\"\n ConvTranspose2d followed by Mish, add, Hardtanh and scaling.\n The expensive convolution is delegated to the same optimized backend as\n PyTorch, while the four following pointwise operations are fused into a\n single custom CUDA kernel that operates in-place on the convolution output.\n \"\"\"\n def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, add_value, scale):\n super(ModelNew, self).__init__()\n self.conv_transpose = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride, padding, output_padding)\n self.add_value = add_value\n self.scale = scale\n\n def forward(self, x):\n y = self.conv_transpose(x)\n return _ext.fused_mish_add_hardtanh_scale(y, float(self.add_value), float(self.scale))\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 7.09 μs
Runtime stats: mean=7.09μs median=?μs std=0.69μs
Numerical precision: max_abs_err=8.34e-06 mean_abs_err=5.27e-07 max_rel_err=1.14e-05
Memory: 5153.42 MB peak (ref: 5153.42 MB, ratio: 1.0x)
Kernel launches: 7 (ref: 16, fusion ratio: 2.2857)
Energy: 2176.36 mJ/run (ref: 1451.04 mJ, ratio: 0.6667x)
SOL score: 1.0
final submitted kernel (89 lines, 3,182 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline
_cpp_src = r"""
#include <torch/extension.h>
torch::Tensor fused_mish_add_hardtanh_scale(torch::Tensor x, double add_value, double scale);
"""
_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>
__device__ __forceinline__ float softplus_stable(float v) {
// Matches the usual stable softplus formulation used by PyTorch closely
// enough for fp32 tolerance, while avoiding overflow for large positives.
if (v > 20.0f) return v;
if (v < -20.0f) return __expf(v);
return log1pf(__expf(v));
}
__device__ __forceinline__ float op_one(float v, float addv, float scale) {
// mish(v) = v * tanh(softplus(v))
float sp = softplus_stable(v);
float y = v * tanhf(sp);
y += addv;
// hardtanh with min=-1, max=1, then scale
y = fminf(1.0f, fmaxf(-1.0f, y));
return y * scale;
}
__global__ void fused_mish_add_hardtanh_scale_kernel(float* __restrict__ data,
int64_t n,
float addv,
float scale) {
int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
int64_t stride = (int64_t)blockDim.x * gridDim.x;
for (; idx < n; idx += stride) {
data[idx] = op_one(data[idx], addv, scale);
}
}
torch::Tensor fused_mish_add_hardtanh_scale(torch::Tensor x, double add_value, double scale) {
const int64_t n = x.numel();
if (n == 0) return x;
const int threads = 256;
int blocks = (int)((n + threads - 1) / threads);
// Keep the grid reasonably sized; the grid-stride loop covers the rest.
if (blocks > 65535) blocks = 65535;
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
fused_mish_add_hardtanh_scale_kernel<<<blocks, threads, 0, stream>>>(
x.data_ptr<float>(), n, (float)add_value, (float)scale);
return x;
}
"""
_ext = load_inline(
name="fused_mish_add_hardtanh_scale_ext",
cpp_sources=_cpp_src,
cuda_sources=_cuda_src,
functions=["fused_mish_add_hardtanh_scale"],
extra_cflags=["-O3"],
extra_cuda_cflags=["-O3", "--use_fast_math"],
verbose=False,
)
class ModelNew(nn.Module):
"""
ConvTranspose2d followed by Mish, add, Hardtanh and scaling.
The expensive convolution is delegated to the same optimized backend as
PyTorch, while the four following pointwise operations are fused into a
single custom CUDA kernel that operates in-place on the convolution output.
"""
def __init__(self, in_channels, out_channels, kernel_size, stride, padding, output_padding, add_value, scale):
super(ModelNew, self).__init__()
self.conv_transpose = nn.ConvTranspose2d(in_channels, out_channels, kernel_size, stride, padding, output_padding)
self.add_value = add_value
self.scale = scale
def forward(self, x):
y = self.conv_transpose(x)
return _ext.fused_mish_add_hardtanh_scale(y, float(self.add_value), float(self.scale))