← Back to Leaderboard

The AI CUDA Engineer 👷

27_SELU_ • 27_selu_manual_unroll_base

Level 1 • Task 27
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 27_selu_manual_unroll_base
CUDA Speedup (Native) 0.942x
CUDA Speedup (Compile) 4.254x
CUDA Runtime 0.007 ms
PyTorch Runtime (Native) 0.007 ms
PyTorch Runtime (Compile) 0.030 ms
Correct True
Max Diff (vs. Reference) 0.000000
Model o3-mini-2025-01-31
Temperature 1.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 and double.
template <typename scalar_t>
__device__ inline scalar_t my_exp(scalar_t x);

template <>
__device__ inline float my_exp<float>(float x) {
    return expf(x);
}

template <>
__device__ inline double my_exp<double>(double x) {
    return exp(x);
}

// CUDA kernel applying SELU activation with manual loop unrolling.
// Each thread processes several elements per iteration to reduce loop overhead.

template <typename scalar_t>
__global__ void selu_kernel_manual_unroll(const scalar_t* __restrict__ input,
                                           scalar_t* __restrict__ output,
                                           size_t numel) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    const int unroll_factor = 4;
    size_t i = tid;

    // Unrolled loop: process unroll_factor elements per iteration
    for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
        #pragma unroll
        for (int j = 0; j < unroll_factor; j++) {
            size_t index = i + j * stride;
            scalar_t x = input[index];
            scalar_t res = (x > static_cast<scalar_t>(0))
                               ? x
                               : static_cast<scalar_t>(1.67326324235437728481) * (my_exp(x) - static_cast<scalar_t>(1));
            output[index] = static_cast<scalar_t>(1.05070098735548049342) * res;
        }
    }

    // Process any remaining elements that don't fit into a full unrolled iteration
    for (; i < numel; i += stride) {
        scalar_t x = input[i];
        scalar_t res = (x > static_cast<scalar_t>(0))
                           ? x
                           : static_cast<scalar_t>(1.67326324235437728481) * (my_exp(x) - static_cast<scalar_t>(1));
        output[i] = static_cast<scalar_t>(1.05070098735548049342) * res;
    }
}

// Host function that launches the manually unrolled CUDA SELU kernel
// Grid dimensions are computed to cover all elements considering the unroll factor.

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

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

    AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "selu_forward_cuda_manual_unroll", ([&] {
        const scalar_t* input_ptr = input.data_ptr<scalar_t>();
        scalar_t* output_ptr = output.data_ptr<scalar_t>();
        selu_kernel_manual_unroll<scalar_t><<<blocks, threads>>>(input_ptr, output_ptr, numel);
    }));

    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &selu_forward, "SELU Activation Forward with Manual Loop Unrolling (CUDA)");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 1.018 inst/cycle 0.000 5
Executed Ipc Elapsed 0.222 inst/cycle 0.000 5
Issue Slots Busy 27.878 % 0.160 5
Issued Ipc Active 1.114 inst/cycle 0.000 5
SM Busy 27.878 % 0.160 5
Memory Throughput 275183327183.384 byte/second 32895306495184756736.000 5
Mem Busy 13.048 % 0.065 5
Max Bandwidth 12.078 % 0.058 5
L1/TEX Hit Rate 0.000 % 0.000 5
L2 Hit Rate 66.982 % 0.019 5
Mem Pipes Busy 2.938 % 0.003 5
Warp Cycles Per Issued Instruction 26.058 cycle 0.007 5
Warp Cycles Per Executed Instruction 28.482 cycle 0.009 5
Avg. Active Threads Per Warp 32.000 0.000 5
Avg. Not Predicated Off Threads Per Warp 28.280 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 2.000 block 0.000 5
Block Limit Shared Mem 8.000 block 0.000 5
Block Limit Warps 2.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 46.198 % 0.014 5
Achieved Active Warps Per SM 29.568 warp 0.006 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 (46.1%) 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 522164.76 μs
Device Time 40.06 μs
Self CPU Time 34.02 μ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 522130.74 μs
Device Time 40.06 μs
Self CPU Time 83.75 μ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 537862.38 μs
Device Time 0.00 μs
Self CPU Time 16170.26 μ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 517900.62 μs
Device Time 0.00 μs
Self CPU Time 517900.62 μ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 400399.92 μs
Device Time 18304.48 μs
Self CPU Time 400399.92 μs
Self Device Time 18304.48 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
void selu_kernel_manual_unroll<float>(float const*, float*, unsigned long)
CPU Time 0.00 μs
Device Time 20467.15 μs
Self CPU Time 0.00 μs
Self Device Time 20467.15 μ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 18031.15 μs
Device Time 35383.46 μs
Self CPU Time 18031.15 μs
Self Device Time 35383.46 μ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 55283.93 μs
Device Time 525873.26 μs
Self CPU Time 9582.06 μ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 45702.76 μs
Device Time 525873.26 μs
Self CPU Time 12447.55 μs
Self Device Time 525873.26 μ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 525951.15 μs
Self CPU Time 0.00 μs
Self Device Time 525951.15 μ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 
45285 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/b2_s3_27_selu_manual_unroll/base/base.cu:27:15 bugprone-narrowing-conversions
27 | int tid = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:28:18: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
28 | int stride = blockDim.x * gridDim.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:33:16: warning: performing an implicit widening conversion to type 'size_t' (aka 'unsigned long') of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
33 | for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:33:16: note: make conversion explicit to silence this warning
5 | for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
| static_cast<size_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:33:16: note: perform multiplication in a wider type
33 | for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
| ^~~~~~~~~~~~~~~~~~
| static_cast<long>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:33:59: warning: performing an implicit widening conversion to type 'size_t' (aka 'unsigned long') of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
33 | for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:33:59: note: make conversion explicit to silence this warning
33 | for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
| ^~~~~~~~~~~~~~~~~~~~~~
| static_cast<size_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:33:59: note: perform multiplication in a wider type
33 | for (; i + (unroll_factor - 1) * stride < numel; i += stride * unroll_factor) {
| ^~~~~~
| static_cast<long>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:36:32: warning: performing an implicit widening conversion to type 'size_t' (aka 'unsigned long') of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
36 | size_t index = i + j * stride;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:36:32: note: make conversion explicit to silence this warning
36 | size_t index = i + j * stride;
| ^~~~~~~~~~
| static_cast<size_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:36:32: note: perform multiplication in a wider type
36 | size_t index = i + j * stride;
| ^
| static_cast<long>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:24: warning: narrowing conversion from 'size_t' (aka 'unsigned long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:33: warning: performing an implicit widening conversion to type 'size_t' (aka 'unsigned long') of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:33: note: make conversion explicit to silence this warning
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^~~~~~~~~~~~~~~~~~~~~~~
| static_cast<size_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:33: note: perform multiplication in a wider type
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^~~~~~~
| static_cast<long>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:65: warning: performing an implicit widening conversion to type 'size_t' (aka 'unsigned long') of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:65: note: make conversion explicit to silence this warning
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^~~~~~~~~~~~~~~~~~~~~~~
| static_cast<size_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:65:65: note: perform multiplication in a wider type
65 | const int blocks = (numel + threads * unroll_factor - 1) / (threads * unroll_factor);
| ^~~~~~~
| static_cast<long>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250202_optimize_b10_s4_e0_sweep/level_1/task_27/b2_s3_27_selu_manual_unroll/base/base.cu:67:5: warning: inside a lambda, '__func__' expands to the name of the function call operator; consider capturing the name of the enclosing function explicitly [bugprone-lambda-function-name]
67 | AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "selu_forward_cuda_manual_unroll", ([&] {
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:237:34: note: expanded from macro 'AT_DISPATCH_FLOATING_TYPES'
237 | AT_DISPATCH_SWITCH(TYPE, NAME, AT_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__))
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:233:3: note: expanded from macro 'AT_DISPATCH_CASE_FLOATING_TYPES'
233 | AT_DISPATCH_CASE(at::ScalarType::Double, __VA_ARGS__) \
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:74:3: note: expanded from macro 'AT_DISPATCH_CASE'
74 | AT_PRIVATE_CASE_TYPE_USING_HINT(enum_type, scalar_t, __VA_ARGS__)
| ^
note: (skipping 1 expansions in backtrace; use -fmacro-backtrace-limit=0 to see all)
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/ATen/Dispatch.h:58:7: note: expanded from macro 'AT_PRIVATE_CHECK_SELECTIVE_BUILD'
58 | AT_ERROR( \
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/c10/util/Exception.h:711:32: note: expanded from macro 'AT_ERROR'
711 | C10_EXPAND_MSVC_WORKAROUND(TORCH_CHECK(false, ::c10::str(__VA_ARGS__))); \
| ^
/home/robert_sakana_ai/miniconda3/envs/llm2cuda/lib/python3.11/site-packages/torch/include/c10/util/Exception.h:536:9: note: expanded from macro 'TORCH_CHECK'
536 | __func__, \
| ^