Kernel Details - selu_vectorized_base

import torch
import torch.nn as nn
import torch.nn.functional as F


def module_fn(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 F.selu(x)


class Model(nn.Module):
    """
    Simple model that performs a SELU activation.
    """

    def __init__(self):
        super(Model, self).__init__()

    def forward(self, x: torch.Tensor, fn=module_fn) -> torch.Tensor:
        return fn(x)


batch_size = 16
dim = 16384


def get_inputs():
    x = torch.randn(batch_size, dim)
    return [x]


def get_init_inputs():
    return []  # No special initialization inputs needed

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 = 16
dim = 16384

def get_inputs():
    x = torch.randn(batch_size, dim)
    return [x]

def get_init_inputs():
    return []  # No special initialization inputs needed

Download Evaluation Download PyTorch Download CUDA Download Profiles

Kernel Information

Operation Name	27_SELU_
Level ID	1
Task ID	27
Kernel Name	selu_vectorized_base_base
CUDA Speedup (Native)	1.099x
CUDA Speedup (Compile)	4.963x
CUDA Runtime	0.006 ms
PyTorch Runtime (Native)	0.007 ms
PyTorch Runtime (Compile)	0.030 ms
Correct	True
Max Diff (vs. Reference)	0.000000
Model	bedrock/anthropic.claude-3-5-sonnet-20241022-v2:0
Temperature	0.00

View Experiment Progress Details

Related Kernels (Level 1, Task 27 • 27_SELU_)

Rank	Kernel Name	Runtime (ms)	Speedup Native	Speedup Compile
🥇	selu_vectorized_base_base	0.01	1.10	4.96
🥇	selu_shared_opt_base	0.01	1.10	4.96
🥇	27_selu_aligned_ldg_base	0.01	1.10	4.96
🥇	27_selu_aligned_ldg_edit_1	0.01	1.10	4.96
5	27_selu_unroll_optimized_base	0.01	0.94	4.25
5	27_SELU_	0.01	0.94	4.25
5	selu_kernel_combined_optimized_base	0.01	0.94	4.25
5	selu_atomic_optimized_base	0.01	0.94	4.25
5	27_selu_manual_unroll_base	0.01	0.94	4.25
5	modular_selu_optimized_base	0.01	0.94	4.25
5	selu_2d_indexing_base	0.01	0.94	4.25
5	selu_shared_mem_optimized_base	0.01	0.94	4.25
5	selu_kernel_combined_base	0.01	0.94	4.25
5	evenly_distributed_selu_base	0.01	0.94	4.25
5	selu_memory_coalesced_base_base	0.01	0.94	4.25
5	evenly_partitioned_selu_base	0.01	0.94	4.25
5	selu_even_load_balance_base	0.01	0.94	4.25
5	selu_atomic_minimal_base	0.01	0.94	4.25
5	selu_coalesced_access_base_base	0.01	0.94	4.25
5	27_selu_vectorized_base	0.01	0.94	4.25

#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>

// Device helper: define an inline exponential function for float
__device__ inline float my_exp(float x) {
    return expf(x);
}

__device__ inline void process_element(float x, float& result) {
    result = (x > 0.0f)
        ? x
        : 1.67326324235437728481f * (my_exp(x) - 1.0f);
    result *= 1.05070098735548049342f;
}

__global__ void selu_kernel_vectorized(const float* __restrict__ input,
                                      float* __restrict__ output,
                                      size_t numel) {
    const size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    const size_t stride = blockDim.x * gridDim.x;
    const size_t vector_stride = stride * 4;
    size_t vector_idx = idx * 4;

    // Process elements in chunks of 4
    for (; vector_idx < (numel & ~3); vector_idx += vector_stride) {
        float4 in_vec = reinterpret_cast<const float4*>(input)[vector_idx >> 2];
        float4 out_vec;

        process_element(in_vec.x, out_vec.x);
        process_element(in_vec.y, out_vec.y);
        process_element(in_vec.z, out_vec.z);
        process_element(in_vec.w, out_vec.w);

        reinterpret_cast<float4*>(output)[vector_idx >> 2] = out_vec;
    }

    // Handle remaining elements
    const size_t remaining_start = numel & ~3;
    for (size_t i = remaining_start + idx; i < numel; i += stride) {
        float result;
        process_element(input[i], result);
        output[i] = result;
    }
}

torch::Tensor selu_forward(torch::Tensor input) {
    TORCH_CHECK(input.is_cuda(), "Input tensor must be a CUDA tensor");
    TORCH_CHECK(input.scalar_type() == torch::kFloat, "Input must be float32");

    auto output = torch::empty_like(input);
    const size_t numel = input.numel();
    const int threads = 256;
    const int blocks = (numel + threads * 4 - 1) / (threads * 4);

    const float* input_ptr = input.data_ptr<float>();
    float* output_ptr = output.data_ptr<float>();
    
    selu_kernel_vectorized<<<blocks, threads>>>(input_ptr, output_ptr, numel);

    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &selu_forward, "SELU Activation Forward with Vectorized Access (CUDA)");
}

Performance Metrics

Metric	Value	Unit	Variance	Samples
Executed Ipc Active	0.450	inst/cycle	0.001	5
Executed Ipc Elapsed	0.180	inst/cycle	0.000	5
Issue Slots Busy	12.756	%	0.463	5
Issued Ipc Active	0.510	inst/cycle	0.001	5
SM Busy	12.756	%	0.463	5
Memory Throughput	287487780320.368	byte/second	1510361079904111104.000	5
Mem Busy	13.698	%	0.001	5
Max Bandwidth	12.644	%	0.001	5
L1/TEX Hit Rate	0.000	%	0.000	5
L2 Hit Rate	67.066	%	0.031	5
Mem Pipes Busy	2.162	%	0.000	5
Warp Cycles Per Issued Instruction	27.052	cycle	2.840	5
Warp Cycles Per Executed Instruction	30.726	cycle	3.663	5
Avg. Active Threads Per Warp	32.000		0.000	5
Avg. Not Predicated Off Threads Per Warp	27.150		0.000	5
Max Active Clusters	0.000	cluster	0.000	5
Max Cluster Size	8.000	block	0.000	5
Overall GPU Occupancy	0.000	%	0.000	5
Cluster Occupancy	0.000	%	0.000	5
Block Limit SM	32.000	block	0.000	5
Block Limit Registers	10.000	block	0.000	5
Block Limit Shared Mem	32.000	block	0.000	5
Block Limit Warps	8.000	block	0.000	5
Theoretical Active Warps per SM	64.000	warp	0.000	5
Theoretical Occupancy	100.000	%	0.000	5
Achieved Occupancy	21.614	%	0.030	5
Achieved Active Warps Per SM	13.834	warp	0.012	5

Analysis Rules

Rule	Description
WRN HighPipeUtilization	All compute pipelines are under-utilized. Either this kernel is very small or it doesn't issue enough warps per scheduler. Check the Launch Statistics and Scheduler Statistics sections for further details.
INF CPIStall	Check the Warp Stall Sampling (All Cycles) table for the top stall locations in your source based on sampling data. The Kernel Profiling Guide (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-reference) provides more details on each stall reason.
WRN Occupancy	This kernel's theoretical occupancy is not impacted by any block limit. The difference between calculated theoretical (100.0%) and measured achieved occupancy (21.5%) can be the result of warp scheduling overheads or workload imbalances during the kernel execution. Load imbalances can occur between warps within a block as well as across blocks of the same kernel. See the CUDA Best Practices Guide (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#occupancy) for more details on optimizing occupancy.

Operation / Metric	Value	Unit
aten::to
CPU Time	739630.06	μs
Device Time	39.97	μs
Self CPU Time	36.54	μs
Self Device Time	0.00	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
aten::_to_copy
CPU Time	739593.52	μs
Device Time	39.97	μs
Self CPU Time	89.55	μs
Self Device Time	0.00	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
aten::empty_strided
CPU Time	758081.83	μs
Device Time	0.00	μs
Self CPU Time	18952.32	μs
Self Device Time	0.00	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
cudaDeviceGetStreamPriorityRange
CPU Time	737894.40	μs
Device Time	0.00	μs
Self CPU Time	737894.40	μs
Self Device Time	0.00	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
cudaLaunchKernel
CPU Time	478002.82	μs
Device Time	21767.01	μs
Self CPU Time	478002.82	μs
Self Device Time	21767.01	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
selu_kernel_vectorized(float const, float, unsigned long)
CPU Time	0.00	μs
Device Time	24149.07	μs
Self CPU Time	0.00	μs
Self Device Time	24149.07	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
cudaEventRecord
CPU Time	17440.72	μs
Device Time	41891.58	μs
Self CPU Time	17440.72	μs
Self Device Time	41891.58	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
aten::zero_
CPU Time	62211.29	μs
Device Time	620076.58	μs
Self CPU Time	12373.77	μs
Self Device Time	0.00	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
aten::fill_
CPU Time	49839.38	μs
Device Time	620076.58	μs
Self CPU Time	16500.91	μs
Self Device Time	620076.58	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B
void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, at::detail::Array<char, 1> >(int, at::native::FillFunctor<int>, at::detail::Array<char, 1>)
CPU Time	0.00	μs
Device Time	620076.58	μs
Self CPU Time	0.00	μs
Self Device Time	620076.58	μs
CPU Memory Usage	0	B
Device Memory Usage	0	B
Self CPU Memory Usage	0	B
Self Device Memory Usage	0	B

Status: Completed

45279 warnings generated when compiling for host.
Suppressed 45321 warnings (45274 in non-user code, 47 NOLINT).
Use -header-filter=.* to display errors from all non-system headers. Use -system-headers to display errors from system headers as well.

/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b6_s3_selu_vectorized_base/base/base.cu:22:27 bugprone-implicit-widening-of-multiplication-result

22 | const size_t stride = blockDim.x * gridDim.x;

| ^

/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b6_s3_selu_vectorized_base/base/base.cu:22:27: note: make conversion explicit to silence this warning

5 | const size_t stride = blockDim.x * gridDim.x;

| ^~~~~~~~~~~~~~~~~~~~~~

| static_cast<const size_t>( )

/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b6_s3_selu_vectorized_base/base/base.cu:22:27: note: perform multiplication in a wider type

22 | const size_t stride = blockDim.x * gridDim.x;

| ^~~~~~~~~~

| static_cast<const size_t>( )

/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b6_s3_selu_vectorized_base/base/base.cu:48:42: warning: the parameter 'input' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]

48 | torch::Tensor selu_forward(torch::Tensor input) {

| ^

| const &

/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b6_s3_selu_vectorized_base/base/base.cu:55:24: warning: narrowing conversion from 'size_t' (aka 'unsigned long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]