← Back to Leaderboard

The AI CUDA Engineer 👷

21_Sigmoidnondivergent_vectorized_sigmoid_base

Level 1 • Task 21
import torch
import torch.nn as nn
import torch.nn.functional as F


def module_fn(x: torch.Tensor) -> torch.Tensor:
    """
    Applies Sigmoid activation to the input tensor.

    Args:
        x (torch.Tensor): Input tensor of any shape.

    Returns:
        torch.Tensor: Output tensor with Sigmoid applied, same shape as input.
    """
    return torch.sigmoid(x)


class Model(nn.Module):
    """
    Simple model that performs a Sigmoid 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 Sigmoid activation.
    """
    def __init__(self):
        super(Model, self).__init__()
    
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        """
        Applies Sigmoid activation to the input tensor.

        Args:
            x (torch.Tensor): Input tensor of any shape.

        Returns:
            torch.Tensor: Output tensor with Sigmoid applied, same shape as input.
        """
        return torch.sigmoid(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

Kernel Information

Related Kernels (Level 1, Task 21 • 21_Sigmoid)

Rank Kernel Name Runtime (ms) Speedup Native Speedup Compile
🥇 sigmoid_shared_mem_optimized_base 0.01 1.11 4.82
🥇 21_sigmoid_modular_device_base 0.01 1.11 4.82
🥇 sigmoid_unroll_optimized_base_base 0.01 1.11 4.82
🥇 sigmoid_min_sync_base_base 0.01 1.11 4.82
🥇 optimized_sigmoid_cuda_base 0.01 1.11 4.82
🥇 optimized_sigmoid_limited_sync_base 0.01 1.11 4.82
🥇 sigmoid_ldg_vectorized_base 0.01 1.11 4.82
🥇 optimized_sigmoid_cuda_base 0.01 1.11 4.82
🥇 optimized_sigmoid_vectorized_combined_edit_1 0.01 1.11 4.82
🥇 21_Sigmoid_optimized_memory_base 0.01 1.11 4.82
🥇 sigmoid_minimal_sync_base_base 0.01 1.11 4.82
🥇 vectorized_ldg_aligned_edit_1 0.01 1.11 4.82
🥇 nondivergent_vectorized_sigmoid_base 0.01 1.11 4.82
🥇 vectorized_no_sync_base 0.01 1.11 4.82
🥇 vectorized_sigmoid_base 0.01 1.11 4.82
🥇 syncthreads_minimal_sigmoid_base 0.01 1.11 4.82
🥇 vectorized_ldg_aligned_base 0.01 1.11 4.82
🥇 optimized_sigmoid_vectorized_combined_base 0.01 1.11 4.82
🥇 optimized_sigmoid_blocksize_tuning_edit_1 0.01 1.11 4.82
🥇 optimized_sigmoid_blocksize_tuning_base 0.01 1.11 4.82
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>

// Vectorized kernel for float using float4 with uniform control flow
__global__ void vectorized_sigmoid_kernel_float(const float* __restrict__ input,
                                                  float* __restrict__ output,
                                                  const int n_vec) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    // Each thread processes a block of consecutive float4 groups
    for (int i = idx; i < n_vec; i += stride) {
        // Load a group of 4 floats at once
        float4 in_val = reinterpret_cast<const float4*>(input)[i];
        float4 out_val;
        // Compute sigmoid uniformly for each component
        out_val.x = 1.0f / (1.0f + expf(-in_val.x));
        out_val.y = 1.0f / (1.0f + expf(-in_val.y));
        out_val.z = 1.0f / (1.0f + expf(-in_val.z));
        out_val.w = 1.0f / (1.0f + expf(-in_val.w));
        reinterpret_cast<float4*>(output)[i] = out_val;
    }
}

// Tail kernel for float to process leftover elements with uniform thread count
__global__ void tail_sigmoid_kernel_float(const float* __restrict__ input,
                                            float* __restrict__ output,
                                            const int start,
                                            const int tail_size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // Launch exactly 'tail_size' threads to minimize divergence in the tail
    if (idx < tail_size) {
        int i = start + idx;
        float in_val = input[i];
        float out_val = 1.0f / (1.0f + expf(-in_val));
        output[i] = out_val;
    }
}

// Fallback scalar kernel for non-float types
template <typename scalar_t>
__global__ void sigmoid_kernel_scalar(const scalar_t* __restrict__ input,
                                        scalar_t* __restrict__ output,
                                        const int64_t size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = idx; i < size; i += stride) {
        float in_val = static_cast<float>(input[i]);
        float out_val = 1.0f / (1.0f + expf(-in_val));
        output[i] = static_cast<scalar_t>(out_val);
    }
}

// Forward function dispatches to vectorized kernels for float and a scalar kernel for others
torch::Tensor forward(torch::Tensor input) {
    auto output = torch::empty_like(input);
    const int64_t size = input.numel();
    const int threads = 256;

    if (input.scalar_type() == at::ScalarType::Float) {
        // Compute the number of groups of 4 elements
        int n_vec = size / 4;
        int tail = size - (n_vec * 4);
        
        // Launch vectorized kernel if there is a complete float4 block
        if (n_vec > 0) {
            int blocks = (n_vec + threads - 1) / threads;
            vectorized_sigmoid_kernel_float<<<blocks, threads>>>(
                input.data_ptr<float>(),
                output.data_ptr<float>(),
                n_vec
            );
        }
        // Launch a separate kernel to handle leftover elements
        if (tail > 0) {
            int blocks_tail = (tail + threads - 1) / threads;
            tail_sigmoid_kernel_float<<<blocks_tail, threads>>>(
                input.data_ptr<float>(),
                output.data_ptr<float>(),
                n_vec * 4,
                tail
            );
        }
    } else {
        int blocks = (size + threads - 1) / threads;
        AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "sigmoid_kernel_scalar", ([&] {
            sigmoid_kernel_scalar<scalar_t><<<blocks, threads>>>(
                input.data_ptr<scalar_t>(),
                output.data_ptr<scalar_t>(),
                size
            );
        }));
    }
    return output;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &forward, "Nondivergent Vectorized Sigmoid forward (CUDA)");
}
Performance Metrics
Metric Value Unit Variance Samples
Analysis Rules
Rule Description
Operation / Metric Value Unit
aten::to
CPU Time 547032.69 μs
Device Time 40.13 μs
Self CPU Time 42.03 μ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 546990.66 μs
Device Time 40.13 μs
Self CPU Time 108.87 μ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 553938.07 μs
Device Time 0.00 μs
Self CPU Time 7434.70 μ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 546269.09 μs
Device Time 0.00 μs
Self CPU Time 546269.09 μ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 149193.17 μs
Device Time 6839.77 μs
Self CPU Time 149193.17 μs
Self Device Time 6839.77 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
vectorized_sigmoid_kernel_float(float const*, float*, int)
CPU Time 0.00 μs
Device Time 7743.70 μs
Self CPU Time 0.00 μs
Self Device Time 7743.70 μ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 7230.24 μs
Device Time 13126.69 μs
Self CPU Time 7230.24 μs
Self Device Time 13126.69 μ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 37944.06 μs
Device Time 202842.31 μs
Self CPU Time 4132.99 μ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 33812.89 μs
Device Time 202842.31 μs
Self CPU Time 4931.00 μs
Self Device Time 202842.31 μ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 202842.31 μs
Self CPU Time 0.00 μs
Self Device Time 202842.31 μ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
45287 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/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:10:15 bugprone-narrowing-conversions
10 | int idx = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:11:18: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
11 | int stride = gridDim.x * blockDim.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:29:45: warning: 2 adjacent parameters of 'tail_sigmoid_kernel_float' of similar type ('const int') are easily swapped by mistake [bugprone-easily-swappable-parameters]
29 | const int start,
| ^~~~~~~~~~~~~~~~
30 | const int tail_size) {
| ~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:29:55: note: the first parameter in the range is 'start'
29 | const int start,
| ^~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:30:55: note: the last parameter in the range is 'tail_size'
30 | const int tail_size) {
| ^~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:31:15: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
31 | int idx = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:46:15: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
46 | int idx = blockIdx.x * blockDim.x + threadIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:47:18: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
47 | int stride = blockDim.x * gridDim.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:63:21: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
63 | int n_vec = size / 4;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:64:20: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
64 | int tail = size - (n_vec * 4);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:64:28: warning: performing an implicit widening conversion to type 'int64_t' (aka 'long') of a multiplication performed in type 'int' [bugprone-implicit-widening-of-multiplication-result]
64 | int tail = size - (n_vec * 4);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:64:28: note: make conversion explicit to silence this warning
5 | int tail = size - (n_vec * 4);
| ^~~~~~~~~
| static_cast<int64_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:64:28: note: perform multiplication in a wider type
64 | int tail = size - (n_vec * 4);
| ^~~~~
| static_cast<int64_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:86:22: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
86 | int blocks = (size + threads - 1) / threads;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_1/task_21/b5_s2_nondivergent_vectorized_sigmoid/base/base.cu:87:9: 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]
87 | AT_DISPATCH_FLOATING_TYPES(input.scalar_type(), "sigmoid_kernel_scalar", ([&] {
| ^
/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__, \
| ^