<commit_before>// ----------------------------------------------------------------
// Gunrock -- Fast and Efficient GPU Graph Library
// ----------------------------------------------------------------
// This source code is distributed under the terms of LICENSE.TXT
// in the root directory of this source distribution.
// ----------------------------------------------------------------


/**
 * @file
 * kernel.cuh
 *
 * @brief Load balanced Edge Map Kernel Entrypoint
 */

#pragma once
#include <gunrock/util/cta_work_distribution.cuh>
#include <gunrock/util/cta_work_progress.cuh>
#include <gunrock/util/kernel_runtime_stats.cuh>

#include <gunrock/oprtr/edge_map_partitioned/cta.cuh>

#include <gunrock/oprtr/advance/kernel_policy.cuh>

namespace gunrock {
namespace oprtr {
namespace edge_map_partitioned {

// GetRowOffsets
//
// RelaxPartitionedEdges

/**
 * Arch dispatch
 */

/**
 * Not valid for this arch (default)
 */
template<
    typename    KernelPolicy,
    typename    ProblemData,
    typename    Functor,
    bool        VALID = (__GR_CUDA_ARCH__ >= KernelPolicy::CUDA_ARCH)>
struct Dispatch
{
    typedef typename KernelPolicy::VertexId VertexId;
    typedef typename KernelPolicy::SizeT    SizeT;
    typedef typename ProblemData::DataSlice DataSlice;

    static __device__ __forceinline__ SizeT GetNeighborListLength(
                            SizeT    *&d_row_offsets,
                            VertexId    *&d_column_indices,
                            VertexId    &d_vertex_id,
                            SizeT       &max_vertex,
                            SizeT       &max_edge,
                            gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
    }

    static __device__ __forceinline__ void GetEdgeCounts(
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_queue,
                                unsigned int *&d_scanned_edges,
                                SizeT &num_elements,
                                SizeT &max_vertex,
                                SizeT &max_edge,
                                gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
    }

    static __device__ __forceinline__ void RelaxPartitionedEdges(
                                bool &queue_reset,
                                VertexId &queue_index,
                                int &label,
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_inverse_column_indices,
                                unsigned int *&d_scanned_edges,
                                unsigned int *&partition_starts,
                                unsigned int &num_partitions,
                                volatile int *&d_done,
                                VertexId *&d_queue,
                                VertexId *&d_out,
                                DataSlice *&problem,
                                SizeT &input_queue_len,
                                SizeT &output_queue_len,
                                SizeT &partition_size,
                                SizeT &max_vertices,
                                SizeT &max_edges,
                                util::CtaWorkProgress &work_progress,
                                util::KernelRuntimeStats &kernel_stats,
                                gunrock::oprtr::advance::TYPE ADVANCE_TYPE,
                                bool &inverse_graph)
    {
    }

    static __device__ __forceinline__ void RelaxLightEdges(
                                bool &queue_reset,
                                VertexId &queue_index,
                                int &label,
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_inverse_column_indices,
                                unsigned int *&d_scanned_edges,
                                volatile int *&d_done,
                                VertexId *&d_queue,
                                VertexId *&d_out,
                                DataSlice *&problem,
                                SizeT &input_queue_len,
                                SizeT &output_queue_len,
                                SizeT &max_vertices,
                                SizeT &max_edges,
                                util::CtaWorkProgress &work_progress,
                                util::KernelRuntimeStats &kernel_stats,
                                gunrock::oprtr::advance::TYPE ADVANCE_TYPE,
                                bool &inverse_graph)
    {
    }

};
template <typename KernelPolicy, typename ProblemData, typename Functor>
struct Dispatch<KernelPolicy, ProblemData, Functor, true>
{
    typedef typename KernelPolicy::VertexId         VertexId;
    typedef typename KernelPolicy::SizeT            SizeT;
    typedef typename ProblemData::DataSlice         DataSlice;

    static __device__ __forceinline__ SizeT GetNeighborListLength(
                            SizeT    *&d_row_offsets,
                            VertexId    *&d_column_indices,
                            VertexId    &d_vertex_id,
                            SizeT       &max_vertex,
                            SizeT       &max_edge,
                            gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
        SizeT first = d_vertex_id >= max_vertex ? max_edge : d_row_offsets[d_vertex_id];
        SizeT second = (d_vertex_id + 1) >= max_vertex ? max_edge : d_row_offsets[d_vertex_id+1];

        return (second > first) ? second - first : 0;
    }

    static __device__ __forceinline__ void GetEdgeCounts(
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_queue,
                                unsigned int *&d_scanned_edges,
                                SizeT &num_elements,
                                SizeT &max_vertex,
                                SizeT &max_edge,
                                gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
        int tid = threadIdx.x;
        int bid = blockIdx.x;

        int my_id = bid*blockDim.x + tid;
        if (my_id >= num_elements || my_id >= max_edge)
            return;
        VertexId v_id = d_queue[my_id];
        SizeT num_edges = GetNeighborListLength(d_row_offsets, d_column_indices, v_id, max_vertex, max_edge, ADVANCE_TYPE);
        d_scanned_edges[my_id] = num_edges;
    }

    static __device__ __forceinline__ void RelaxPartitionedEdges(
                                bool &queue_reset,
                                VertexId &queue_index,
                                int &label,
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_inverse_column_indices,
                                unsigned int *&d_scanned_edges,
                                unsigned int *&partition_starts,
                                unsigned int &num_partitions,<commit_msg>handle -1 vertex id input in load-balanced advance operator<commit_after>// ----------------------------------------------------------------
// Gunrock -- Fast and Efficient GPU Graph Library
// ----------------------------------------------------------------
// This source code is distributed under the terms of LICENSE.TXT
// in the root directory of this source distribution.
// ----------------------------------------------------------------


/**
 * @file
 * kernel.cuh
 *
 * @brief Load balanced Edge Map Kernel Entrypoint
 */

#pragma once
#include <gunrock/util/cta_work_distribution.cuh>
#include <gunrock/util/cta_work_progress.cuh>
#include <gunrock/util/kernel_runtime_stats.cuh>

#include <gunrock/oprtr/edge_map_partitioned/cta.cuh>

#include <gunrock/oprtr/advance/kernel_policy.cuh>

namespace gunrock {
namespace oprtr {
namespace edge_map_partitioned {

// GetRowOffsets
//
// RelaxPartitionedEdges

/**
 * Arch dispatch
 */

/**
 * Not valid for this arch (default)
 */
template<
    typename    KernelPolicy,
    typename    ProblemData,
    typename    Functor,
    bool        VALID = (__GR_CUDA_ARCH__ >= KernelPolicy::CUDA_ARCH)>
struct Dispatch
{
    typedef typename KernelPolicy::VertexId VertexId;
    typedef typename KernelPolicy::SizeT    SizeT;
    typedef typename ProblemData::DataSlice DataSlice;

    static __device__ __forceinline__ SizeT GetNeighborListLength(
                            SizeT    *&d_row_offsets,
                            VertexId    *&d_column_indices,
                            VertexId    &d_vertex_id,
                            SizeT       &max_vertex,
                            SizeT       &max_edge,
                            gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
    }

    static __device__ __forceinline__ void GetEdgeCounts(
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_queue,
                                unsigned int *&d_scanned_edges,
                                SizeT &num_elements,
                                SizeT &max_vertex,
                                SizeT &max_edge,
                                gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
    }

    static __device__ __forceinline__ void RelaxPartitionedEdges(
                                bool &queue_reset,
                                VertexId &queue_index,
                                int &label,
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_inverse_column_indices,
                                unsigned int *&d_scanned_edges,
                                unsigned int *&partition_starts,
                                unsigned int &num_partitions,
                                volatile int *&d_done,
                                VertexId *&d_queue,
                                VertexId *&d_out,
                                DataSlice *&problem,
                                SizeT &input_queue_len,
                                SizeT &output_queue_len,
                                SizeT &partition_size,
                                SizeT &max_vertices,
                                SizeT &max_edges,
                                util::CtaWorkProgress &work_progress,
                                util::KernelRuntimeStats &kernel_stats,
                                gunrock::oprtr::advance::TYPE ADVANCE_TYPE,
                                bool &inverse_graph)
    {
    }

    static __device__ __forceinline__ void RelaxLightEdges(
                                bool &queue_reset,
                                VertexId &queue_index,
                                int &label,
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_inverse_column_indices,
                                unsigned int *&d_scanned_edges,
                                volatile int *&d_done,
                                VertexId *&d_queue,
                                VertexId *&d_out,
                                DataSlice *&problem,
                                SizeT &input_queue_len,
                                SizeT &output_queue_len,
                                SizeT &max_vertices,
                                SizeT &max_edges,
                                util::CtaWorkProgress &work_progress,
                                util::KernelRuntimeStats &kernel_stats,
                                gunrock::oprtr::advance::TYPE ADVANCE_TYPE,
                                bool &inverse_graph)
    {
    }

};
template <typename KernelPolicy, typename ProblemData, typename Functor>
struct Dispatch<KernelPolicy, ProblemData, Functor, true>
{
    typedef typename KernelPolicy::VertexId         VertexId;
    typedef typename KernelPolicy::SizeT            SizeT;
    typedef typename ProblemData::DataSlice         DataSlice;

    static __device__ __forceinline__ SizeT GetNeighborListLength(
                            SizeT    *&d_row_offsets,
                            VertexId    *&d_column_indices,
                            VertexId    &d_vertex_id,
                            SizeT       &max_vertex,
                            SizeT       &max_edge,
                            gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
        SizeT first = d_vertex_id >= max_vertex ? max_edge : d_row_offsets[d_vertex_id];
        SizeT second = (d_vertex_id + 1) >= max_vertex ? max_edge : d_row_offsets[d_vertex_id+1];

        return (second > first) ? second - first : 0;
    }

    static __device__ __forceinline__ void GetEdgeCounts(
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_queue,
                                unsigned int *&d_scanned_edges,
                                SizeT &num_elements,
                                SizeT &max_vertex,
                                SizeT &max_edge,
                                gunrock::oprtr::advance::TYPE &ADVANCE_TYPE)
    {
        int tid = threadIdx.x;
        int bid = blockIdx.x;

        int my_id = bid*blockDim.x + tid;
        if (my_id >= num_elements || my_id >= max_edge)
            return;
        VertexId v_id = d_queue[my_id];
        if (v_id == -1) {
            d_scanned_edges[my_id] = 0;
            return;
        }
        SizeT num_edges = GetNeighborListLength(d_row_offsets, d_column_indices, v_id, max_vertex, max_edge, ADVANCE_TYPE);
        d_scanned_edges[my_id] = num_edges;
    }

    static __device__ __forceinline__ void RelaxPartitionedEdges(
                                bool &queue_reset,
                                VertexId &queue_index,
                                int &label,
                                SizeT *&d_row_offsets,
                                VertexId *&d_column_indices,
                                VertexId *&d_inverse_column_indices,
                                unsigned int *&d_scanned_edges,
                                unsigned int *&partition_starts,
                                unsigned int &num_partitions,
{'fraction_non_alphanumeric': 0.08769063180827887, 'fraction_numerical': 0.0006127450980392157, 'mean_word_length': 1.1147422977253096, 'pattern_counts': {'":': 0, '<': 19, '<?xml version=': 0, '>': 31, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 6, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '6098053', 'n_tokens_mistral': 3841, 'n_tokens_neox': 3590, 'n_words': 737}