correct
27_SELU_.py
started
2026-05-02T02:14:35Z
timingkernelbench.timing
kernel mean / median / std
4.34 / — / 0.08 μs
ref mean / median / std
4.30 / — / 0.09 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
1.00x
energyNVML (pynvml)
ratio (ref / kernel)
0.81x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
3029.81 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):
"""
Simple model that performs a SELU activation.
"""
def __init__(self):
super(Model, self).__init__()
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""
Applies SELU activation to the input tensor.
Args:
x (torch.Tensor): Input tensor of any shape.
Returns:
torch.Tensor: Output tensor with SELU applied, same shape as input.
"""
return torch.selu(x)
batch_size = 4096
dim = 393216
def get_inputs():
x = torch.rand(batch_size, dim)
return [x]
def get_init_inputs():
return [] # No special initialization inputs neededturn 05.2s turn · 5.2s LLM · 0.0s tools
reasoning
Let's analyze the problem:
1. We need to replace `torch.selu(x)` with a custom CUDA kernel.
2. The input shape is `[batch_size, dim] = [4096, 393216]`.
3. This is a large tensor with ~1.6 billion elements (4096 * 393216 = 1,610,612,736).
4. Each element is independent, so this is an element-wise kernel.
5. For large element-wise operations, performance is memory bandwidth bound.
6. We need to write an efficient fused CUDA kernel that loads inputs, applies SELU, and stores outputs.
7. SELU formula: `scale * (max(0, x) + min(0, alpha * (exp(x) - 1)))` where `scale ≈ 1.0507` and `alpha ≈ 1.6733`.
Let me check the exact PyTorch SELU parameters:
- alpha = 1.6732632423543772848170429916717
- scale = 1.0507010221481323242187531591474
The formula is:
```
selu(x) = scale * x if x > 0
selu(x) = scale * alpha * (exp(x) - 1) if x <= 0
```
For a kernel of 1.6B elements, we need to maximize memory throughput. Standard approaches:
- Use vectorized loads/stores (float4)
- Use grid-stride loops
- Ensure coalesced memory access
Let me get the GPU specs first to understand memory bandwidth and optimal block size.
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 1142.7s turn · 142.7s LLM · 0.0s tools
reasoning
The task is to replace `torch.selu` with a custom CUDA kernel for H100. Let me analyze what's needed:
1. Input: [4096, 393216] = 1,610,612,736 elements (~1.6 billion)
2. `torch.selu(x)` is a simple element-wise operation: `scale * max(0, x) + scale * alpha * (exp(min(0, x)) - 1)`
where scale ≈ 1.0507, alpha ≈ 1.67326
For H100 with 3.35 TB/s bandwidth, this is clearly memory-bandwidth bound. The input is fp32 (4 bytes), output is fp32 (4 bytes). Total memory traffic is ~12.88 GB. At 3.35 TB/s, theoretical minimum time is ~3.84 ms.
To maximize performance, I need to:
1. Use a simple CUDA kernel with vectorized loads/stores (e.g., float4)
2. Use enough threads per block (256 or 512)
3. Ensure grid size is large enough to saturate the GPU
Let me write a kernel:
- Process 4 elements per thread using float4
- Grid-stride loop
For SELU constants:
- scale = 1.0507009873554804934193349852946
- alpha = 1.6732632423543772848180429911717
Formula: `scale * x` if x > 0 else `scale * alpha * (exp(x) - 1)`
Let me write the kernel code.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 2271.0s turn · 271.0s LLM · 0.0s tools
reasoning
Good, it compiled. Now let's check correctness.
call → run_correctness
callrun_correctness
{
"kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
turn 3101.1s turn · 101.1s LLM · 0.0s tools
reasoning
The kernel compiles and passes correctness. Let me now run a static analysis before submitting.
call → static_check
callstatic_check
{
"kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n"
}
result ← static_check (ok)
okstatic_check
static_check PASSED: no violations or warnings detected.
turn 429.8s turn · 29.8s LLM · 0.0s tools · FINAL
reasoning
The kernel passes correctness and static checks. Let me now submit it. But before I submit, let me review once more:
1. The kernel uses float4 vectorized loads/stores with proper alignment considerations.
2. The SELU formula is correctly implemented with the right constants.
3. The kernel handles the tail case where n is not a multiple of 4.
4. The block size is 256 which is good.
5. `-O3` and `--use_fast_math` flags are used for optimization.
However, I should double-check: the input `x = torch.rand(batch_size, dim)` will be in the range [0, 1), so all values are positive. This means the SELU will always take the positive branch. Does my kernel handle this correctly? Yes, because `x > 0` evaluates to true, and we return `scale * x`.
One concern: `float4` requires 16-byte alignment. For tensors allocated by PyTorch, this should generally be satisfied since CUDA memory allocations are typically aligned to at least 256 bytes. And `batch_size * dim = 4096 * 393216 = 1,610,612,736`, which is divisible by 4 (and by 16), so the total number of elements is a multiple of 4.
Actually, wait. The starting address of the tensor data should be aligned to at least 16 bytes for float4. PyTorch CUDA tensors are allocated with 256-byte alignment typically, so this should be fine.
Let me submit the kernel.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "\nimport torch\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\nconstexpr float SELU_SCALE = 1.0507009873554804934193349852946f;\nconstexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;\n\n__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {\n constexpr int VEC_SIZE = 4;\n int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;\n \n if (idx + VEC_SIZE <= n) {\n float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];\n float4 out;\n \n out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);\n out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);\n out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);\n out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);\n \n reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;\n } else {\n for (int64_t i = idx; i < n; ++i) {\n float x = input[i];\n output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);\n }\n }\n}\n\ntorch::Tensor selu_forward(torch::Tensor input) {\n auto output = torch::empty_like(input);\n int64_t n = input.numel();\n \n if (n == 0) return output;\n \n constexpr int VEC_SIZE = 4;\n constexpr int BLOCK_SIZE = 256;\n int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;\n int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;\n \n selu_kernel<<<num_blocks, BLOCK_SIZE>>>(\n input.data_ptr<float>(),\n output.data_ptr<float>(),\n n\n );\n \n return output;\n}\n'''\n\nselu_module = load_inline(\n name='selu_cuda',\n cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',\n cuda_sources=cuda_source,\n functions=['selu_forward'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False\n)\n\nclass ModelNew(torch.nn.Module):\n def __init__(self):\n super(ModelNew, self).__init__()\n \n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return selu_module.selu_forward(x)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 4.34 μs
Runtime stats: mean=4.34μs median=?μs std=0.08μs
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
Memory: 12288.0 MB peak (ref: 12288.0 MB, ratio: 1.0x)
Kernel launches: 2 (ref: 2, fusion ratio: 1.0)
Energy: 2417.6 mJ/run (ref: 1953.18 mJ, ratio: 0.8079x)
SOL score: 0.9044
final submitted kernel (70 lines, 2,251 chars) — full source, untruncated
import torch
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#include <cuda_runtime.h>
constexpr float SELU_SCALE = 1.0507009873554804934193349852946f;
constexpr float SELU_ALPHA = 1.6732632423543772848180429911717f;
__global__ void selu_kernel(const float* __restrict__ input, float* __restrict__ output, int64_t n) {
constexpr int VEC_SIZE = 4;
int64_t idx = (blockIdx.x * blockDim.x + threadIdx.x) * VEC_SIZE;
if (idx + VEC_SIZE <= n) {
float4 in = reinterpret_cast<const float4*>(input)[idx / VEC_SIZE];
float4 out;
out.x = in.x > 0.0f ? SELU_SCALE * in.x : SELU_SCALE * SELU_ALPHA * (expf(in.x) - 1.0f);
out.y = in.y > 0.0f ? SELU_SCALE * in.y : SELU_SCALE * SELU_ALPHA * (expf(in.y) - 1.0f);
out.z = in.z > 0.0f ? SELU_SCALE * in.z : SELU_SCALE * SELU_ALPHA * (expf(in.z) - 1.0f);
out.w = in.w > 0.0f ? SELU_SCALE * in.w : SELU_SCALE * SELU_ALPHA * (expf(in.w) - 1.0f);
reinterpret_cast<float4*>(output)[idx / VEC_SIZE] = out;
} else {
for (int64_t i = idx; i < n; ++i) {
float x = input[i];
output[i] = x > 0.0f ? SELU_SCALE * x : SELU_SCALE * SELU_ALPHA * (expf(x) - 1.0f);
}
}
}
torch::Tensor selu_forward(torch::Tensor input) {
auto output = torch::empty_like(input);
int64_t n = input.numel();
if (n == 0) return output;
constexpr int VEC_SIZE = 4;
constexpr int BLOCK_SIZE = 256;
int64_t num_vec = (n + VEC_SIZE - 1) / VEC_SIZE;
int64_t num_blocks = (num_vec + BLOCK_SIZE - 1) / BLOCK_SIZE;
selu_kernel<<<num_blocks, BLOCK_SIZE>>>(
input.data_ptr<float>(),
output.data_ptr<float>(),
n
);
return output;
}
'''
selu_module = load_inline(
name='selu_cuda',
cpp_sources='torch::Tensor selu_forward(torch::Tensor input);',
cuda_sources=cuda_source,
functions=['selu_forward'],
extra_cuda_cflags=['-O3', '--use_fast_math'],
verbose=False
)
class ModelNew(torch.nn.Module):
def __init__(self):
super(ModelNew, self).__init__()
def forward(self, x: torch.Tensor) -> torch.Tensor:
return selu_module.selu_forward(x)