auto_sort_bitonic.hpp

Go to the documentation of this file.

// Copyright (c) 2021 - 2023 SparseLinearAlgebra
// Autogenerated file, do not modify
 
#pragma once
 
static const char source_sort_bitonic[] = R"(
 
 
 
// memory bank conflict-free address and local buffer size
#ifdef LM_NUM_MEM_BANKS
    #define LM_ADDR(address) (address + ((address) / LM_NUM_MEM_BANKS))
    #define LM_SIZE(size)    (size + (size) / LM_NUM_MEM_BANKS)
#endif
 
#define SWAP_KEYS(x, y) \
    uint tmp1 = x;      \
    x         = y;      \
    y         = tmp1;
 
#define SWAP_VALUES(x, y) \
    TYPE tmp2 = x;        \
    x         = y;        \
    y         = tmp2;
 
// nearest power of two number greater equals n
uint ceil_to_pow2(uint n) {
    uint r = 1;
    while (r < n) r *= 2;
    return r;
}
 
// find first element in a sorted array such x <= element
uint lower_bound(const uint           x,
                 uint                 first,
                 uint                 size,
                 __global const uint* array) {
    while (size > 0) {
        int step = size / 2;
 
        if (array[first + step] < x) {
            first = first + step + 1;
            size -= step + 1;
        } else {
            size = step;
        }
    }
    return first;
}
 
// find first element in a sorted array such x <= element
uint lower_bound_local(const uint          x,
                       uint                first,
                       uint                size,
                       __local const uint* array) {
    while (size > 0) {
        int step = size / 2;
 
        if (array[first + step] < x) {
            first = first + step + 1;
            size -= step + 1;
        } else {
            size = step;
        }
    }
    return first;
}
__kernel void bitonic_sort_local(__global uint* g_keys,
                                 __global TYPE* g_values,
                                 const uint     total_n) {
    const uint grid  = get_group_id(0);
    const uint lid   = get_local_id(0);
    const uint lsize = get_local_size(0);
 
    const uint offset    = grid * BLOCK_SIZE;
    const uint border    = min(offset + BLOCK_SIZE, total_n);
    const uint n         = border - offset;
    const uint n_aligned = ceil_to_pow2(n);
    const uint n_threads = n_aligned / 2;
 
    __local uint s_keys[BLOCK_SIZE];
    __local TYPE s_values[BLOCK_SIZE];
 
    for (uint i = lid; i + offset < border; i += lsize) {
        s_keys[i]   = g_keys[i + offset];
        s_values[i] = g_values[i + offset];
    }
 
    barrier(CLK_LOCAL_MEM_FENCE);
 
    for (uint segment_size = 2; segment_size <= n_aligned; segment_size *= 2) {
        const uint segment_size_half = segment_size / 2;
 
        for (uint tid = lid; tid < n_threads; tid += lsize) {
            const uint segment_id       = tid / segment_size_half;
            const uint inner_id         = tid % segment_size_half;
            const uint inner_id_sibling = segment_size - inner_id - 1;
            const uint i                = segment_id * segment_size + inner_id;
            const uint j                = segment_id * segment_size + inner_id_sibling;
 
            if (i < n && j < n && s_keys[i] > s_keys[j]) {
                SWAP_KEYS(s_keys[i], s_keys[j]);
                SWAP_VALUES(s_values[i], s_values[j]);
            }
        }
 
        barrier(CLK_LOCAL_MEM_FENCE);
 
        for (uint k = segment_size_half / 2; k > 0; k /= 2) {
            for (uint tid = lid; tid < n_threads; tid += lsize) {
                const uint segment_size_inner = k * 2;
                const uint segment_id         = tid / k;
                const uint inner_id           = tid % k;
                const uint inner_id_sibling   = inner_id + k;
                const uint i                  = segment_id * segment_size_inner + inner_id;
                const uint j                  = segment_id * segment_size_inner + inner_id_sibling;
 
                if (i < n && j < n && s_keys[i] > s_keys[j]) {
                    SWAP_KEYS(s_keys[i], s_keys[j]);
                    SWAP_VALUES(s_values[i], s_values[j]);
                }
            }
 
            barrier(CLK_LOCAL_MEM_FENCE);
        }
    }
 
    for (uint i = lid; i + offset < border; i += lsize) {
        g_keys[i + offset]   = s_keys[i];
        g_values[i + offset] = s_values[i];
    }
}
 
__kernel void bitonic_sort_global(__global uint* g_keys,
                                  __global TYPE* g_values,
                                  const uint     n,
                                  const uint     segment_start) {
    const uint lid       = get_local_id(0);
    const uint lsize     = get_local_size(0);
    const uint n_aligned = ceil_to_pow2(n);
    const uint n_threads = n_aligned / 2;
 
    for (uint segment_size = segment_start; segment_size <= n_aligned; segment_size *= 2) {
        const uint segment_size_half = segment_size / 2;
 
        for (uint tid = lid; tid < n_threads; tid += lsize) {
            const uint segment_id       = tid / segment_size_half;
            const uint inner_id         = tid % segment_size_half;
            const uint inner_id_sibling = segment_size - inner_id - 1;
            const uint i                = segment_id * segment_size + inner_id;
            const uint j                = segment_id * segment_size + inner_id_sibling;
 
            if (i < n && j < n && g_keys[i] > g_keys[j]) {
                SWAP_KEYS(g_keys[i], g_keys[j]);
                SWAP_VALUES(g_values[i], g_values[j]);
            }
        }
 
        barrier(CLK_GLOBAL_MEM_FENCE);
 
        for (uint k = segment_size_half / 2; k > 0; k /= 2) {
            for (uint tid = lid; tid < n_threads; tid += lsize) {
                const uint segment_size_inner = k * 2;
                const uint segment_id         = tid / k;
                const uint inner_id           = tid % k;
                const uint inner_id_sibling   = inner_id + k;
                const uint i                  = segment_id * segment_size_inner + inner_id;
                const uint j                  = segment_id * segment_size_inner + inner_id_sibling;
 
                if (i < n && j < n && g_keys[i] > g_keys[j]) {
                    SWAP_KEYS(g_keys[i], g_keys[j]);
                    SWAP_VALUES(g_values[i], g_values[j]);
                }
            }
 
            barrier(CLK_GLOBAL_MEM_FENCE);
        }
    }
}
)";