correct
47_Conv3d_Mish_Tanh.py
started
2026-05-02T03:05:39Z
timingkernelbench.timing
kernel mean / median / std
4.13 / — / 0.09 μs
ref mean / median / std
2.56 / — / 0.01 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
1.29x
energyNVML (pynvml)
ratio (ref / kernel)
0.60x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
5932.26 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 3D convolution, applies Mish activation, and then applies Tanh activation.
"""
def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):
super(Model, self).__init__()
self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)
def forward(self, x):
"""
Args:
x (torch.Tensor): Input tensor of shape (batch_size, in_channels, D, H, W).
Returns:
torch.Tensor: Output tensor of shape (batch_size, out_channels, D', H', W').
"""
x = self.conv(x)
x = torch.nn.functional.mish(x)
x = torch.tanh(x)
return x
batch_size = 16
in_channels = 32
out_channels = 64
D, H, W = 32, 64, 64
kernel_size = 3
def get_inputs():
return [torch.rand(batch_size, in_channels, D, H, W)]
def get_init_inputs():
return [in_channels, out_channels, kernel_size]turn 015.6s turn · 15.6s LLM · 0.0s tools
call → get_gpu_specs
callget_gpu_specs
{}
result ← get_gpu_specs (ok)
okget_gpu_specs
GPU specs for NVIDIA H100 80GB HBM3:
total memory (runtime): 79.1 GB
spec entry: H100
GPU Architecture: Hopper
GPU Memory: 80GB
Memory Bandwidth: 3.35 TB/s
FP64 TFLOPS: 34
FP64 Tensor Core TFLOPS: 67
FP32 TFLOPS: 67
TF32 Tensor Core TFLOPS: 989 with sparsity
BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity
FP16 Tensor Core TFLOPS: 1979 with sparsity
FP8 Tensor Core TFLOPS: 3958 with sparsity
INT8 Tensor Core TOPS: 3958 with sparsity
Register File Size: 64K 32-bit registers per SM
Maximum number of registers per thread: 255
Maximum number of thread blocks per SM: 32
Shared memory capacity per SM: 228 KB
Maximum shared memory per thread block: 227 KB
turn 141.9s turn · 41.9s LLM · 0.0s tools
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\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <vector>\n\n__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) {\n int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;\n int64_t stride = (int64_t)blockDim.x * gridDim.x;\n for (int64_t i = idx; i < n; i += stride) {\n float x = data[i];\n // softplus stable enough for conv output range; exact branch avoids overflow\n float sp;\n if (x > 20.0f) sp = x;\n else if (x < -20.0f) sp = expf(x);\n else sp = log1pf(expf(x));\n float mish = x * tanhf(sp);\n data[i] = tanhf(mish);\n }\n}\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding) {\n auto y = torch::conv3d(input, weight, bias, {stride, stride, stride}, {padding, padding, padding}, {1,1,1}, 1);\n const int64_t n = y.numel();\n if (n > 0) {\n int threads = 256;\n int blocks = (int)((n + threads - 1) / threads);\n if (blocks > 4096) blocks = 4096;\n mish_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='conv_mish_tanh_ext_v1', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['conv_mish_tanh'], verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):\n super(ModelNew, self).__init__()\n self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)\n self.stride = stride if isinstance(stride, int) else stride[0]\n self.padding = padding if isinstance(padding, int) else padding[0]\n\n def forward(self, x):\n return _ext.conv_mish_tanh(x, self.conv.weight, self.conv.bias, int(self.stride), int(self.padding))\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 25.0s turn · 5.0s LLM · 0.0s tools
call → run_correctness
callrun_correctness
{
"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\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <vector>\n\n__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) {\n int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;\n int64_t stride = (int64_t)blockDim.x * gridDim.x;\n for (int64_t i = idx; i < n; i += stride) {\n float x = data[i];\n float sp;\n if (x > 20.0f) sp = x;\n else if (x < -20.0f) sp = expf(x);\n else sp = log1pf(expf(x));\n float mish = x * tanhf(sp);\n data[i] = tanhf(mish);\n }\n}\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding) {\n auto y = torch::conv3d(input, weight, bias, {stride, stride, stride}, {padding, padding, padding}, {1,1,1}, 1);\n const int64_t n = y.numel();\n if (n > 0) {\n int threads = 256;\n int blocks = (int)((n + threads - 1) / threads);\n if (blocks > 4096) blocks = 4096;\n mish_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='conv_mish_tanh_ext_v1c', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['conv_mish_tanh'], verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):\n super(ModelNew, self).__init__()\n self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)\n self.stride = stride if isinstance(stride, int) else stride[0]\n self.padding = padding if isinstance(padding, int) else padding[0]\n\n def forward(self, x):\n return _ext.conv_mish_tanh(x, self.conv.weight, self.conv.bias, int(self.stride), int(self.padding))\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=1.02e-05 mean_abs_err=6.89e-07 max_rel_err=1.66e-05
turn 324.0s turn · 24.0s LLM · 0.0s tools
reasoning
**Optimizing convolution timing**
I'm considering how to optimize this process, but the final timing is still unknown. I could submit it, yet I think I need to aim for faster activation. The reference activation might use two kernels, and I'm analyzing the performance from PyTorch.
I’m estimating the outputs and looking into using cuDNN for better efficiency. I need to ensure that my custom implementation aligns with the expected tolerances, especially concerning the tensor cores and operand types. Finding the optimal approach is vital.
**Refining activation functions**
I'm wondering if I should skip using `--use_fast_math` for better accuracy, which seems to pass fine. I can change the formula for Mish to mish=x*tanh(log1p(exp(x)). I think limiting x to around [-1, 1] could help. For convolution, Kaiming weights might have bounds as sqrt values, and I need to consider how many terms I'm dealing with in my calculations.
I see that the bottleneck is in transcendentals instead of memory, so I should explore vectorizing with float4. Plus, maybe fewer calls for Mish and tanh could enhance performance. A polynomial or rational approximation could work!
call → static_check
callstatic_check
{
"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\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <vector>\n\n__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) {\n int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;\n int64_t stride = (int64_t)blockDim.x * gridDim.x;\n for (int64_t i = idx; i < n; i += stride) {\n float x = data[i];\n float sp;\n if (x > 20.0f) sp = x;\n else if (x < -20.0f) sp = expf(x);\n else sp = log1pf(expf(x));\n float mish = x * tanhf(sp);\n data[i] = tanhf(mish);\n }\n}\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding) {\n auto y = torch::conv3d(input, weight, bias, {stride, stride, stride}, {padding, padding, padding}, {1,1,1}, 1);\n const int64_t n = y.numel();\n if (n > 0) {\n int threads = 256;\n int blocks = (int)((n + threads - 1) / threads);\n if (blocks > 4096) blocks = 4096;\n mish_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='conv_mish_tanh_ext_v1s', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['conv_mish_tanh'], verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):\n super(ModelNew, self).__init__()\n self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)\n self.stride = stride if isinstance(stride, int) else stride[0]\n self.padding = padding if isinstance(padding, int) else padding[0]\n\n def forward(self, x):\n return _ext.conv_mish_tanh(x, self.conv.weight, self.conv.bias, int(self.stride), int(self.padding))\n"
}
result ← static_check (ok)
okstatic_check
static_check PASSED: no violations or warnings detected.
turn 427.3s turn · 27.3s LLM · 0.0s tools
reasoning
**Considering code generation**
I need to generate a source for submission that includes code for computing coefficients and embedding. I can use numpy to handle polynomial fitting with Chebyshev, which requires the source to be a string. But if numpy isn’t available, I might use torch's linear algebra tools instead. The kernel code should be generated at module load before compilation, and while testing performance can be tricky, ensuring correctness is essential. There is some risk involved with polynomial methods, but they could potentially offer speed advantages with FMAs.
**Evaluating output ranges**
I'm looking at the outputs for the function defined for values outside the range of [-5,5]. For x > 5, I expect the output to be around 1, and for x < -5, it should be a small negative value. I need to ensure that my understanding of how the tanh function behaves in these regions is correct. Let’s test the implementation for accuracy and verify that it aligns with my expectations!
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport numpy as _np\n\n# Build a high-accuracy Chebyshev approximation for tanh(mish(x)) on the range\n# that contains essentially all Conv3d outputs for the benchmark initialization.\n_xs = _np.linspace(-5.0, 5.0, 300001, dtype=_np.float64)\n_ys = _np.tanh(_xs * _np.tanh(_np.log1p(_np.exp(_xs))))\n_coef = _np.polynomial.chebyshev.chebfit(_xs / 5.0, _ys, 36).astype(_np.float32)\n_lines = []\n_lines.append('float t = x * 0.2f;')\n_lines.append('float b0 = 0.0f, b1 = 0.0f, b2 = 0.0f;')\nfor c in _coef[:0:-1]:\n _lines.append('b2 = b1; b1 = b0; b0 = fmaf(2.0f*t, b1, %.9g); b0 -= b2;' % float(c))\n_lines.append('return fmaf(t, b1, %.9g - b2);' % float(_coef[0]))\n_poly_body = '\\n '.join(_lines)\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <vector>\n\n__device__ __forceinline__ float exact_act(float x) {\n float sp;\n if (x > 20.0f) sp = x;\n else if (x < -20.0f) sp = expf(x);\n else sp = log1pf(expf(x));\n return tanhf(x * tanhf(sp));\n}\n\n__device__ __forceinline__ float poly_act(float x) {\n if (x < -5.0f || x > 5.0f) return exact_act(x);\n ''' + _poly_body + r'''\n}\n\n__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) {\n int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;\n int64_t stride = (int64_t)blockDim.x * gridDim.x;\n for (int64_t i = idx; i < n; i += stride) {\n float x = data[i];\n data[i] = poly_act(x);\n }\n}\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding) {\n auto y = torch::conv3d(input, weight, bias, {stride, stride, stride}, {padding, padding, padding}, {1,1,1}, 1);\n const int64_t n = y.numel();\n if (n > 0) {\n int threads = 256;\n int blocks = (int)((n + threads - 1) / threads);\n if (blocks > 4096) blocks = 4096;\n mish_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='conv_mish_tanh_ext_poly1', cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=['conv_mish_tanh'], verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'], extra_cflags=['-O3'])\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):\n super(ModelNew, self).__init__()\n self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)\n self.stride = stride if isinstance(stride, int) else stride[0]\n self.padding = padding if isinstance(padding, int) else padding[0]\n\n def forward(self, x):\n return _ext.conv_mish_tanh(x, self.conv.weight, self.conv.bias, int(self.stride), int(self.padding))\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: (0 / 5) trials did not all match.
correctness_issue: Output mismatch
max_difference: ['0.921518', '0.921518', '0.921518', '0.921518', '0.921518']
avg_difference: ['0.823098', '0.823105', '0.823117', '0.823094', '0.823096']
turn 510.0s turn · 10.0s LLM · 0.0s tools
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport numpy as _np\n_xs = _np.linspace(-5.0, 5.0, 300001, dtype=_np.float64)\n_ys = _np.tanh(_xs * _np.tanh(_np.log1p(_np.exp(_xs))))\n_coef = _np.polynomial.chebyshev.chebfit(_xs / 5.0, _ys, 36).astype(_np.float32)\n_lines = []\n_lines.append('float t = x * 0.2f;')\n_lines.append('float b0 = 0.0f, b1 = 0.0f, b2 = 0.0f;')\nfor c in _coef[:0:-1]:\n _lines.append('b2 = b1; b1 = b0; b0 = fmaf(2.0f*t, b1, %.9g - b2);' % float(c))\n_lines.append('return fmaf(t, b0, %.9g - b1);' % float(_coef[0]))\n_poly_body = '\\n '.join(_lines)\n_cpp_src = r'''\n#include <torch/extension.h>\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,int64_t stride, int64_t padding);\n'''\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <vector>\n__device__ __forceinline__ float exact_act(float x) { float sp; if (x > 20.0f) sp = x; else if (x < -20.0f) sp = expf(x); else sp = log1pf(expf(x)); return tanhf(x * tanhf(sp)); }\n__device__ __forceinline__ float poly_act(float x) { if (x < -5.0f || x > 5.0f) return exact_act(x);\n ''' + _poly_body + r'''\n}\n__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) { int64_t idx=(int64_t)blockIdx.x*blockDim.x+threadIdx.x; int64_t stride=(int64_t)blockDim.x*gridDim.x; for(int64_t i=idx;i<n;i+=stride){ data[i]=poly_act(data[i]); }}\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,int64_t stride, int64_t padding) { auto y=torch::conv3d(input, weight, bias, {stride,stride,stride}, {padding,padding,padding}, {1,1,1}, 1); const int64_t n=y.numel(); if(n>0){ int threads=256; int blocks=(int)((n+threads-1)/threads); if(blocks>4096) blocks=4096; mish_tanh_kernel<<<blocks,threads,0,at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);} return y;}\n'''\n_ext = load_inline(name='conv_mish_tanh_ext_poly2', cpp_sources=_cpp_src, cuda_sources=_cuda_src, functions=['conv_mish_tanh'], verbose=False, extra_cuda_cflags=['-O3','--use_fast_math'], extra_cflags=['-O3'])\nclass ModelNew(nn.Module):\n def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):\n super(ModelNew,self).__init__(); self.conv=nn.Conv3d(in_channels,out_channels,kernel_size,stride=stride,padding=padding); self.stride=stride if isinstance(stride,int) else stride[0]; self.padding=padding if isinstance(padding,int) else padding[0]\n def forward(self,x):\n return _ext.conv_mish_tanh(x,self.conv.weight,self.conv.bias,int(self.stride),int(self.padding))\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=4.01e-05 mean_abs_err=2.24e-05 max_rel_err=2.54e+07
turn 68.9s turn · 8.9s LLM · 0.0s tools
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport numpy as _np\n_deg=24\n_xs = _np.linspace(-5.0, 5.0, 200001, dtype=_np.float64)\n_ys = _np.tanh(_xs * _np.tanh(_np.log1p(_np.exp(_xs))))\n_coef = _np.polynomial.chebyshev.chebfit(_xs / 5.0, _ys, _deg).astype(_np.float32)\n_lines=['float t = x * 0.2f;','float b0 = 0.0f, b1 = 0.0f, b2 = 0.0f;']\nfor c in _coef[:0:-1]: _lines.append('b2 = b1; b1 = b0; b0 = fmaf(2.0f*t, b1, %.9g - b2);' % float(c))\n_lines.append('return fmaf(t, b0, %.9g - b1);' % float(_coef[0]))\n_poly_body='\\n '.join(_lines)\n_cpp_src='#include <torch/extension.h>\\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,int64_t stride, int64_t padding);\\n'\n_cuda_src=r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n__device__ __forceinline__ float exact_act(float x){float sp;if(x>20.0f)sp=x;else if(x<-20.0f)sp=expf(x);else sp=log1pf(expf(x));return tanhf(x*tanhf(sp));}\n__device__ __forceinline__ float poly_act(float x){if(x<-5.0f||x>5.0f)return exact_act(x);\n '''+_poly_body+r'''\n}\n__global__ void mish_tanh_kernel(float* __restrict__ data,int64_t n){int64_t idx=(int64_t)blockIdx.x*blockDim.x+threadIdx.x;int64_t stride=(int64_t)blockDim.x*gridDim.x;for(int64_t i=idx;i<n;i+=stride)data[i]=poly_act(data[i]);}\ntorch::Tensor conv_mish_tanh(torch::Tensor input,torch::Tensor weight,torch::Tensor bias,int64_t stride,int64_t padding){auto y=torch::conv3d(input,weight,bias,{stride,stride,stride},{padding,padding,padding},{1,1,1},1);int64_t n=y.numel();if(n){int threads=256;int blocks=(int)((n+threads-1)/threads);if(blocks>4096)blocks=4096;mish_tanh_kernel<<<blocks,threads,0,at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(),n);}return y;}\n'''\n_ext=load_inline(name='conv_mish_tanh_ext_poly24',cpp_sources=_cpp_src,cuda_sources=_cuda_src,functions=['conv_mish_tanh'],verbose=False,extra_cuda_cflags=['-O3','--use_fast_math'],extra_cflags=['-O3'])\nclass ModelNew(nn.Module):\n def __init__(self,in_channels,out_channels,kernel_size,stride=1,padding=0):\n super(ModelNew,self).__init__();self.conv=nn.Conv3d(in_channels,out_channels,kernel_size,stride=stride,padding=padding);self.stride=stride if isinstance(stride,int) else stride[0];self.padding=padding if isinstance(padding,int) else padding[0]\n def forward(self,x): return _ext.conv_mish_tanh(x,self.conv.weight,self.conv.bias,int(self.stride),int(self.padding))\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: (0 / 5) trials did not all match.
correctness_issue: Output mismatch
max_difference: ['0.000940', '0.000940', '0.000940', '0.000940', '0.000940']
avg_difference: ['0.000576', '0.000576', '0.000576', '0.000576', '0.000576']
turn 77.8s turn · 7.8s LLM · 0.0s tools
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport numpy as _np\n_deg=32\n_xs=_np.linspace(-5,5,250001,dtype=_np.float64); _ys=_np.tanh(_xs*_np.tanh(_np.log1p(_np.exp(_xs)))); _coef=_np.polynomial.chebyshev.chebfit(_xs/5,_ys,_deg).astype(_np.float32)\n_lines=['float t = x * 0.2f;','float b0 = 0.0f, b1 = 0.0f, b2 = 0.0f;']\nfor c in _coef[:0:-1]: _lines.append('b2 = b1; b1 = b0; b0 = fmaf(2.0f*t, b1, %.9g - b2);'%float(c))\n_lines.append('return fmaf(t, b0, %.9g - b1);'%float(_coef[0])); _poly='\\n'.join(_lines)\n_cpp_src='#include <torch/extension.h>\\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,int64_t stride, int64_t padding);\\n'\n_cuda_src='''#include <torch/extension.h>\\n#include <ATen/cuda/CUDAContext.h>\\n__device__ __forceinline__ float exact_act(float x){float sp;if(x>20.0f)sp=x;else if(x<-20.0f)sp=expf(x);else sp=log1pf(expf(x));return tanhf(x*tanhf(sp));}\\n__device__ __forceinline__ float poly_act(float x){if(x<-5.0f||x>5.0f)return exact_act(x);\\n'''+_poly+'''\\n}\\n__global__ void mish_tanh_kernel(float* __restrict__ data,int64_t n){int64_t idx=(int64_t)blockIdx.x*blockDim.x+threadIdx.x;int64_t stride=(int64_t)blockDim.x*gridDim.x;for(int64_t i=idx;i<n;i+=stride)data[i]=poly_act(data[i]);}\\ntorch::Tensor conv_mish_tanh(torch::Tensor input,torch::Tensor weight,torch::Tensor bias,int64_t stride,int64_t padding){auto y=torch::conv3d(input,weight,bias,{stride,stride,stride},{padding,padding,padding},{1,1,1},1);int64_t n=y.numel();if(n){int threads=256;int blocks=(int)((n+threads-1)/threads);if(blocks>4096)blocks=4096;mish_tanh_kernel<<<blocks,threads,0,at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(),n);}return y;}\\n'''\n_ext=load_inline(name='conv_mish_tanh_ext_poly32',cpp_sources=_cpp_src,cuda_sources=_cuda_src,functions=['conv_mish_tanh'],verbose=False,extra_cuda_cflags=['-O3','--use_fast_math'],extra_cflags=['-O3'])\nclass ModelNew(nn.Module):\n def __init__(self,in_channels,out_channels,kernel_size,stride=1,padding=0):\n super(ModelNew,self).__init__(); self.conv=nn.Conv3d(in_channels,out_channels,kernel_size,stride=stride,padding=padding); self.stride=stride if isinstance(stride,int) else stride[0]; self.padding=padding if isinstance(padding,int) else padding[0]\n def forward(self,x): return _ext.conv_mish_tanh(x,self.conv.weight,self.conv.bias,int(self.stride),int(self.padding))\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: (0 / 5) trials did not all match.
correctness_issue: Output mismatch
max_difference: ['0.000122', '0.000122', '0.000122', '0.000122', '0.000122']
avg_difference: ['0.000071', '0.000071', '0.000071', '0.000071', '0.000071']
turn 818.6s turn · 18.6s LLM · 0.0s tools · FINAL
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\nimport numpy as _np\n\n# Chebyshev approximation coefficients are generated once at module load and\n# embedded into the CUDA source before compilation. This replaces the two\n# expensive activation kernels with a single lightweight in-place kernel after\n# cuDNN's convolution.\n_xs = _np.linspace(-5.0, 5.0, 300001, dtype=_np.float64)\n_ys = _np.tanh(_xs * _np.tanh(_np.log1p(_np.exp(_xs))))\n_coef = _np.polynomial.chebyshev.chebfit(_xs / 5.0, _ys, 36).astype(_np.float32)\n_lines = ['float t = x * 0.2f;', 'float b0 = 0.0f, b1 = 0.0f, b2 = 0.0f;']\nfor _c in _coef[:0:-1]:\n _lines.append('b2 = b1; b1 = b0; b0 = fmaf(2.0f*t, b1, %.9g - b2);' % float(_c))\n_lines.append('return fmaf(t, b0, %.9g - b1);' % float(_coef[0]))\n_poly_body = '\\n '.join(_lines)\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n\n__device__ __forceinline__ float exact_act(float x) {\n float sp;\n if (x > 20.0f) sp = x;\n else if (x < -20.0f) sp = expf(x);\n else sp = log1pf(expf(x));\n return tanhf(x * tanhf(sp));\n}\n\n__device__ __forceinline__ float poly_act(float x) {\n if (x < -5.0f || x > 5.0f) return exact_act(x);\n ''' + _poly_body + r'''\n}\n\n__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) {\n int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;\n int64_t grid_stride = (int64_t)blockDim.x * gridDim.x;\n for (int64_t i = idx; i < n; i += grid_stride) {\n data[i] = poly_act(data[i]);\n }\n}\n\ntorch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,\n int64_t stride, int64_t padding) {\n auto y = torch::conv3d(input, weight, bias,\n {stride, stride, stride},\n {padding, padding, padding},\n {1, 1, 1}, 1);\n const int64_t n = y.numel();\n if (n > 0) {\n const int threads = 256;\n int blocks = (int)((n + threads - 1) / threads);\n if (blocks > 4096) blocks = 4096;\n mish_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='conv_mish_tanh_ext_final',\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=['conv_mish_tanh'],\n verbose=False,\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n extra_cflags=['-O3'])\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):\n super(ModelNew, self).__init__()\n self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)\n self.stride = stride if isinstance(stride, int) else stride[0]\n self.padding = padding if isinstance(padding, int) else padding[0]\n\n def forward(self, x):\n return _ext.conv_mish_tanh(x, self.conv.weight, self.conv.bias, int(self.stride), int(self.padding))\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 4.13 μs
Runtime stats: mean=4.13μs median=?μs std=0.09μs
Numerical precision: max_abs_err=4.01e-05 mean_abs_err=2.24e-05 max_rel_err=2.54e+07
Memory: 963.1 MB peak (ref: 1157.36 MB, ratio: 0.8322x)
Kernel launches: 7 (ref: 9, fusion ratio: 1.2857)
Energy: 2194.68 mJ/run (ref: 1311.02 mJ, ratio: 0.5974x)
SOL score: 1.0
final submitted kernel (85 lines, 3,304 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
import numpy as _np
# Chebyshev approximation coefficients are generated once at module load and
# embedded into the CUDA source before compilation. This replaces the two
# expensive activation kernels with a single lightweight in-place kernel after
# cuDNN's convolution.
_xs = _np.linspace(-5.0, 5.0, 300001, dtype=_np.float64)
_ys = _np.tanh(_xs * _np.tanh(_np.log1p(_np.exp(_xs))))
_coef = _np.polynomial.chebyshev.chebfit(_xs / 5.0, _ys, 36).astype(_np.float32)
_lines = ['float t = x * 0.2f;', 'float b0 = 0.0f, b1 = 0.0f, b2 = 0.0f;']
for _c in _coef[:0:-1]:
_lines.append('b2 = b1; b1 = b0; b0 = fmaf(2.0f*t, b1, %.9g - b2);' % float(_c))
_lines.append('return fmaf(t, b0, %.9g - b1);' % float(_coef[0]))
_poly_body = '\n '.join(_lines)
_cpp_src = r'''
#include <torch/extension.h>
torch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,
int64_t stride, int64_t padding);
'''
_cuda_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
__device__ __forceinline__ float exact_act(float x) {
float sp;
if (x > 20.0f) sp = x;
else if (x < -20.0f) sp = expf(x);
else sp = log1pf(expf(x));
return tanhf(x * tanhf(sp));
}
__device__ __forceinline__ float poly_act(float x) {
if (x < -5.0f || x > 5.0f) return exact_act(x);
''' + _poly_body + r'''
}
__global__ void mish_tanh_kernel(float* __restrict__ data, int64_t n) {
int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
int64_t grid_stride = (int64_t)blockDim.x * gridDim.x;
for (int64_t i = idx; i < n; i += grid_stride) {
data[i] = poly_act(data[i]);
}
}
torch::Tensor conv_mish_tanh(torch::Tensor input, torch::Tensor weight, torch::Tensor bias,
int64_t stride, int64_t padding) {
auto y = torch::conv3d(input, weight, bias,
{stride, stride, stride},
{padding, padding, padding},
{1, 1, 1}, 1);
const int64_t n = y.numel();
if (n > 0) {
const int threads = 256;
int blocks = (int)((n + threads - 1) / threads);
if (blocks > 4096) blocks = 4096;
mish_tanh_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(y.data_ptr<float>(), n);
}
return y;
}
'''
_ext = load_inline(name='conv_mish_tanh_ext_final',
cpp_sources=_cpp_src,
cuda_sources=_cuda_src,
functions=['conv_mish_tanh'],
verbose=False,
extra_cuda_cflags=['-O3', '--use_fast_math'],
extra_cflags=['-O3'])
class ModelNew(nn.Module):
def __init__(self, in_channels, out_channels, kernel_size, stride=1, padding=0):
super(ModelNew, self).__init__()
self.conv = nn.Conv3d(in_channels, out_channels, kernel_size, stride=stride, padding=padding)
self.stride = stride if isinstance(stride, int) else stride[0]
self.padding = padding if isinstance(padding, int) else padding[0]
def forward(self, x):
return _ext.conv_mish_tanh(x, self.conv.weight, self.conv.bias, int(self.stride), int(self.padding))