auto_reduce_by_key.hpp

Go to the documentation of this file.

 
// Copyright (c) 2021 - 2023 SparseLinearAlgebra
// Autogenerated file, do not modify
 
#pragma once
 
static const char source_reduce_by_key[] = 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;
}
// generate uint offsets for unique keys to store result
__kernel void reduce_by_key_generate_offsets(__global const uint* g_keys,
                                             __global uint*       g_offsets,
                                             const uint           n) {
    const uint gid = get_global_id(0);
 
    if (gid < n) {
        bool is_neq    = gid + 1 < n && g_keys[gid] != g_keys[gid + 1];
        g_offsets[gid] = is_neq ? 1 : 0;
    }
}
 
// scalar reduction for each group of keys
__kernel void reduce_by_key_scalar(__global const uint* g_keys,
                                   __global const TYPE* g_values,
                                   __global const uint* g_offsets,
                                   __global uint*       g_unique_keys,
                                   __global TYPE*       g_reduce_values,
                                   const uint           n_keys,
                                   const uint           n_groups) {
    const uint gid = get_global_id(0);
 
    if (gid < n_groups) {
        const uint start_idx = lower_bound(gid, 0, n_keys, g_offsets);
        TYPE       value     = g_values[start_idx];
 
        for (uint i = start_idx + 1; i < n_keys && gid == g_offsets[i]; i += 1) {
            value = OP_BINARY(value, g_values[i]);
        }
 
        g_unique_keys[gid]   = g_keys[start_idx];
        g_reduce_values[gid] = value;
    }
}
 
__kernel void reduce_by_key_small(__global const uint* g_keys,
                                  __global const TYPE* g_values,
                                  __global uint*       g_unique_keys,
                                  __global TYPE*       g_reduce_values,
                                  __global uint*       g_reduced_count,
                                  const uint           n_keys) {
    const uint lid       = get_local_id(0);
    const uint n_aligned = ceil_to_pow2(n_keys);
 
    __local uint s_offsets[BLOCK_SIZE];
 
    uint gen_key = 0;
    if (lid < n_keys) {
        bool is_neq   = lid > 0 && g_keys[lid] != g_keys[lid - 1];
        bool is_first = lid == 0;
        gen_key       = is_neq || is_first ? 1 : 0;
    }
    s_offsets[lid] = gen_key;
 
    for (uint offset = 1; offset < n_aligned; offset *= 2) {
        barrier(CLK_LOCAL_MEM_FENCE);
        uint value = s_offsets[lid];
 
        if (offset <= lid) {
            value += s_offsets[lid - offset];
        }
 
        barrier(CLK_LOCAL_MEM_FENCE);
        s_offsets[lid] = value;
    }
 
    barrier(CLK_LOCAL_MEM_FENCE);
    const uint n_values = s_offsets[n_keys - 1];
 
    if (lid < n_values) {
        const uint id        = lid + 1;
        const uint start_idx = lower_bound_local(id, 0, n_keys, s_offsets);
        TYPE       value     = g_values[start_idx];
 
        for (uint i = start_idx + 1; i < n_keys && id == s_offsets[i]; i += 1) {
            value = OP_BINARY(value, g_values[i]);
        }
 
        g_unique_keys[lid]   = g_keys[start_idx];
        g_reduce_values[lid] = value;
    }
 
    if (lid == 0) {
        g_reduced_count[0] = n_values;
    }
}
 
__kernel void reduce_by_key_sequential(__global const uint* g_keys,
                                       __global const TYPE* g_values,
                                       __global uint*       g_unique_keys,
                                       __global TYPE*       g_reduce_values,
                                       __global uint*       g_reduced_count,
                                       const uint           n_keys) {
    const uint gid = get_global_id(0);
 
    if (gid == 0) {
        uint count         = 0;
        uint current_key   = g_keys[0];
        TYPE current_value = g_values[0];
 
        for (uint read_offset = 1; read_offset < n_keys; read_offset += 1) {
            if (g_keys[read_offset] == current_key) {
                current_value = OP_BINARY(current_value, g_values[read_offset]);
            } else {
                g_unique_keys[count]   = current_key;
                g_reduce_values[count] = current_value;
                current_key            = g_keys[read_offset];
                current_value          = g_values[read_offset];
                count += 1;
            }
        }
 
        g_unique_keys[count]   = current_key;
        g_reduce_values[count] = current_value;
        g_reduced_count[0]     = count + 1;
    }
}
)";