<commit_before>/******************************************************************************
 * Copyright (c) 2011, <NAME>.  All rights reserved.
 * Copyright (c) 2011-2014, NVIDIA CORPORATION.  All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *     * Redistributions of source code must retain the above copyright
 *       notice, this list of conditions and the following disclaimer.
 *     * Redistributions in binary form must reproduce the above copyright
 *       notice, this list of conditions and the following disclaimer in the
 *       documentation and/or other materials provided with the distribution.
 *     * Neither the name of the NVIDIA CORPORATION nor the
 *       names of its contributors may be used to endorse or promote products
 *       derived from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 ******************************************************************************/

/******************************************************************************
 * Test of BlockReduce utilities
 ******************************************************************************/

// Ensure printing of CUDA runtime errors to console
#define CUB_STDERR

#include <stdio.h>

#include <device_functions.h>

#include <cub/util_ptx.cuh>
#include <cub/util_allocator.cuh>
#include <cub/block/block_load.cuh>
#include <cub/block/block_reduce.cuh>

#include "test_util.h"

using namespace cub;


//---------------------------------------------------------------------
// Globals, constants and typedefs
//---------------------------------------------------------------------

bool                    g_verbose       = false;
int                     g_repeat        = 0;
CachingDeviceAllocator  g_allocator(true);



//---------------------------------------------------------------------
// Test kernels
//---------------------------------------------------------------------


/// Generic reduction (full, 1)
template <typename BlockReduce, typename T, typename ReductionOp>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T (&data)[1], ReductionOp &reduction_op)
{
    return BlockReduce(temp_storage).Reduce(data[0], reduction_op);
}

/// Generic reduction (full, ITEMS_PER_THREAD)
template <typename BlockReduce, typename T, int ITEMS_PER_THREAD, typename ReductionOp>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T (&data)[ITEMS_PER_THREAD], ReductionOp &reduction_op)
{
    return BlockReduce(temp_storage).Reduce(data, reduction_op);
}

/// Generic reduction (partial, 1)
template <typename BlockReduce, typename T, typename ReductionOp>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T &data, ReductionOp &reduction_op, int valid_threads)
{
    return BlockReduce(temp_storage).Reduce(data, reduction_op, valid_threads);
}

/// Sum reduction (full, 1)
template <typename BlockReduce, typename T>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T (&data)[1], Sum &reduction_op)
{
    return BlockReduce(temp_storage).Sum(data[0]);
}

/// Sum reduction (full, ITEMS_PER_THREAD)
template <typename BlockReduce, typename T, int ITEMS_PER_THREAD>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T (&data)[ITEMS_PER_THREAD], Sum &reduction_op)
{
    return BlockReduce(temp_storage).Sum(data);
}

/// Sum reduction (partial, 1)
template <typename BlockReduce, typename T>
__device__ __forceinline__ T DeviceTest(
    typename BlockReduce::TempStorage &temp_storage, T &data, Sum &reduction_op, int valid_threads)
{
    return BlockReduce(temp_storage).Sum(data, valid_threads);
}


/**
 * Test full-tile reduction kernel (where num_items is an even
 * multiple of BLOCK_THREADS)
 */
template <
    BlockReduceAlgorithm    ALGORITHM,
    int                     BLOCK_DIM_X,
    int                     BLOCK_DIM_Y,
    int                     BLOCK_DIM_Z,
    int                     ITEMS_PER_THREAD,
    typename                T,
    typename                ReductionOp>
__launch_bounds__ (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)
__global__ void FullTileReduceKernel(
    T                       *d_in,
    T                       *d_out,
    ReductionOp             reduction_op,
    int                     tiles,
    clock_t                 *d_elapsed)
{
    const int BLOCK_THREADS     = BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z;
    const int TILE_SIZE         = BLOCK_THREADS * ITEMS_PER_THREAD;

    // Cooperative threadblock reduction utility type (returns aggregate in thread 0)
    typedef BlockReduce<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockReduce;

    // Allocate temp storage in shared memory
    __shared__ typename BlockReduce::TempStorage temp_storage;

    int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);

    // Per-thread tile data
    T data[ITEMS_PER_THREAD];

    // Load first tile of data
    int block_offset = 0;
    LoadDirectBlocked(linear_tid, d_in + block_offset, data);
    block_offset += TILE_SIZE;

    // Start cycle timer
    clock_t start = clock();

    // Cooperative reduce first tile
    T block_aggregate = DeviceTest<BlockReduce>(temp_storage, data, reduction_op);

    // Stop cycle timer
#if CUB_PTX_ARCH == 100
    // Bug: recording stop clock causes mis-write of running prefix value
    clock_t stop = 0;
#else
    clock_t stop = clock();
#endif // CUB_PTX_ARCH == 100
    clock_t elapsed = (start > stop) ? start - stop : stop - start;

    // Loop over input tiles
    while (block_offset < TILE_SIZE * tiles)
    {
        // TestBarrier between threadblock reductions
        __syncthreads();

        // Load tile of data
        LoadDirectBlocked(linear_tid, d_in + block_offset, data);
        block_offset += TILE_SIZE;

        // Start cycle timer
        clock_t start = clock();

        // Cooperatively reduce the tile's aggregate
        T tile_aggregate = DeviceTest<BlockReduce>(temp_storage, data, reduction_op);

        // Stop cycle timer
    #if CUB_PTX_ARCH == 100
        // Bug: recording stop clock causes mis-write of running prefix value
        clock_t stop = 0;
    #else
        clock_t stop = clock();
    #endif // CUB_PTX_ARCH == 100
        elapsed += (start > stop) ? start - stop : stop - start;

        // Reduce threadblock aggregate
        block_aggregate = reduction_op(block_aggregate, tile_aggregate);
    }

    // Store data
    if (linear_tid == 0)
    {
        d_out[0] = block_aggregate;
        *d_elapsed = elapsed;
    }
}



/**
 * Test partial-tile reduction kernel (where num_items < BLOCK_THREADS)
 */
template <
    BlockReduceAlgorithm    ALGORITHM,
    int                     BLOCK_DIM_X,
    int                     BLOCK_DIM_Y,
    int                     BLOCK_DIM_Z,
    typename                T,
    typename                ReductionOp>
__launch_bounds__ (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)
__global__ void PartialTileReduceKernel(
    T                       *d_in,
    T                       *d_out,
    int                     num_items,
    ReductionOp             reduction_op,
    clock_t                 *d_elapsed)
{
    // Cooperative threadblock reduction utility type (returns aggregate only in thread-0)
    typedef BlockReduce<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockReduce;

    // Allocate temp storage in shared memory
    __shared__ typename BlockReduce::TempStorage temp_storage;

    int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);

    // Per-thread tile data
    T partial;

    // Load partial tile data
    if (linear_tid < num_items)
    {
        partial = d_in[linear_tid];
    }

    // Start cycle timer
    clock_t start = clock();

    // Cooperatively reduce the tile's aggregate
    T tile_aggregate = DeviceTest<BlockReduce>(temp_storage, partial, reduction_op, num_items);

    // Stop cycle timer
#if CUB_PTX_ARCH == 100
    // Bug: recording stop clock causes mis-write of running prefix value
    clock_t stop = 0;
#else
    clock_t stop = clock();
#endif // CUB_PTX_ARCH == 100

    clock_t elapsed = (start > stop) ? start - stop : stop - start;
<commit_msg>Fix block reduce test<commit_after>/******************************************************************************
 * Copyright (c) 2011, <NAME>.  All rights reserved.
 * Copyright (c) 2011-2014, NVIDIA CORPORATION.  All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *     * Redistributions of source code must retain the above copyright
 *       notice, this list of conditions and the following disclaimer.
 *     * Redistributions in binary form must reproduce the above copyright
 *       notice, this list of conditions and the following disclaimer in the
 *       documentation and/or other materials provided with the distribution.
 *     * Neither the name of the NVIDIA CORPORATION nor the
 *       names of its contributors may be used to endorse or promote products
 *       derived from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 ******************************************************************************/

/******************************************************************************
 * Test of BlockReduce utilities
 ******************************************************************************/

// Ensure printing of CUDA runtime errors to console
#define CUB_STDERR

#include <stdio.h>

#include <device_functions.h>

#include <cub/util_ptx.cuh>
#include <cub/util_allocator.cuh>
#include <cub/block/block_load.cuh>
#include <cub/block/block_reduce.cuh>

#include "test_util.h"

using namespace cub;


//---------------------------------------------------------------------
// Globals, constants and typedefs
//---------------------------------------------------------------------

bool                    g_verbose       = false;
int                     g_repeat        = 0;
CachingDeviceAllocator  g_allocator(true);



//---------------------------------------------------------------------
// Test kernels
//---------------------------------------------------------------------


/// Generic reduction (full, 1)
template <typename BlockReduceT, typename T, typename ReductionOp>
__device__ __forceinline__ T DeviceTest(
    BlockReduceT &block_reduce, T (&data)[1], ReductionOp &reduction_op)
{
    return block_reduce.Reduce(data[0], reduction_op);
}

/// Generic reduction (full, ITEMS_PER_THREAD)
template <typename BlockReduceT, typename T, int ITEMS_PER_THREAD, typename ReductionOp>
__device__ __forceinline__ T DeviceTest(
    BlockReduceT &block_reduce, T (&data)[ITEMS_PER_THREAD], ReductionOp &reduction_op)
{
    return block_reduce.Reduce(data, reduction_op);
}

/// Generic reduction (partial, 1)
template <typename BlockReduceT, typename T, typename ReductionOp>
__device__ __forceinline__ T DeviceTest(
    BlockReduceT &block_reduce, T &data, ReductionOp &reduction_op, int valid_threads)
{
    return block_reduce.Reduce(data, reduction_op, valid_threads);
}

/// Sum reduction (full, 1)
template <typename BlockReduceT, typename T>
__device__ __forceinline__ T DeviceTest(
    BlockReduceT &block_reduce, T (&data)[1], Sum &reduction_op)
{
    return block_reduce.Sum(data[0]);
}

/// Sum reduction (full, ITEMS_PER_THREAD)
template <typename BlockReduceT, typename T, int ITEMS_PER_THREAD>
__device__ __forceinline__ T DeviceTest(
    BlockReduceT &block_reduce, T (&data)[ITEMS_PER_THREAD], Sum &reduction_op)
{
    return block_reduce.Sum(data);
}

/// Sum reduction (partial, 1)
template <typename BlockReduceT, typename T>
__device__ __forceinline__ T DeviceTest(
    BlockReduceT &block_reduce, T &data, Sum &reduction_op, int valid_threads)
{
    return block_reduce.Sum(data, valid_threads);
}


/**
 * Test full-tile reduction kernel (where num_items is an even
 * multiple of BLOCK_THREADS)
 */
template <
    BlockReduceAlgorithm    ALGORITHM,
    int                     BLOCK_DIM_X,
    int                     BLOCK_DIM_Y,
    int                     BLOCK_DIM_Z,
    int                     ITEMS_PER_THREAD,
    typename                T,
    typename                ReductionOp>
__launch_bounds__ (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)
__global__ void FullTileReduceKernel(
    T                       *d_in,
    T                       *d_out,
    ReductionOp             reduction_op,
    int                     tiles,
    clock_t                 *d_elapsed)
{
    const int BLOCK_THREADS     = BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z;
    const int TILE_SIZE         = BLOCK_THREADS * ITEMS_PER_THREAD;

    // Cooperative threadblock reduction utility type (returns aggregate in thread 0)
    typedef BlockReduce<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockReduceT;

    // Allocate temp storage in shared memory
    __shared__ typename BlockReduceT::TempStorage temp_storage;

    int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);

    // Per-thread tile data
    T data[ITEMS_PER_THREAD];

    // Load first tile of data
    int block_offset = 0;
    LoadDirectBlocked(linear_tid, d_in + block_offset, data);
    block_offset += TILE_SIZE;

    // Start cycle timer
    clock_t start = clock();

    // Cooperative reduce first tile
    BlockReduceT block_reduce(temp_storage) ;
    T block_aggregate = DeviceTest(block_reduce, data, reduction_op);

    // Stop cycle timer
#if CUB_PTX_ARCH == 100
    // Bug: recording stop clock causes mis-write of running prefix value
    clock_t stop = 0;
#else
    clock_t stop = clock();
#endif // CUB_PTX_ARCH == 100
    clock_t elapsed = (start > stop) ? start - stop : stop - start;

    // Loop over input tiles
    while (block_offset < TILE_SIZE * tiles)
    {
        // TestBarrier between threadblock reductions
        __syncthreads();

        // Load tile of data
        LoadDirectBlocked(linear_tid, d_in + block_offset, data);
        block_offset += TILE_SIZE;

        // Start cycle timer
        clock_t start = clock();

        // Cooperatively reduce the tile's aggregate
        T tile_aggregate = DeviceTest(block_reduce, data, reduction_op);

        // Stop cycle timer
    #if CUB_PTX_ARCH == 100
        // Bug: recording stop clock causes mis-write of running prefix value
        clock_t stop = 0;
    #else
        clock_t stop = clock();
    #endif // CUB_PTX_ARCH == 100
        elapsed += (start > stop) ? start - stop : stop - start;

        // Reduce threadblock aggregate
        block_aggregate = reduction_op(block_aggregate, tile_aggregate);
    }

    // Store data
    if (linear_tid == 0)
    {
        d_out[0] = block_aggregate;
        *d_elapsed = elapsed;
    }
}



/**
 * Test partial-tile reduction kernel (where num_items < BLOCK_THREADS)
 */
template <
    BlockReduceAlgorithm    ALGORITHM,
    int                     BLOCK_DIM_X,
    int                     BLOCK_DIM_Y,
    int                     BLOCK_DIM_Z,
    typename                T,
    typename                ReductionOp>
__launch_bounds__ (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z)
__global__ void PartialTileReduceKernel(
    T                       *d_in,
    T                       *d_out,
    int                     num_items,
    ReductionOp             reduction_op,
    clock_t                 *d_elapsed)
{
    // Cooperative threadblock reduction utility type (returns aggregate only in thread-0)
    typedef BlockReduce<T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z> BlockReduceT;

    // Allocate temp storage in shared memory
    __shared__ typename BlockReduceT::TempStorage temp_storage;

    int linear_tid = RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z);

    // Per-thread tile data
    T partial;

    // Load partial tile data
    if (linear_tid < num_items)
    {
        partial = d_in[linear_tid];
    }

    // Start cycle timer
    clock_t start = clock();

    // Cooperatively reduce the tile's aggregate
    BlockReduceT block_reduce(temp_storage) ;
    T tile_aggregate = DeviceTest(block_reduce, partial, reduction_op, num_items);

    // Stop cycle timer
#if CUB_PTX_ARCH == 100
    // Bug: recording stop clock causes mis-write of running prefix value
    clock_t stop = 0;
#else
    clock_t stop = clock();
#endif // CUB_PTX_ARCH == 100

    clock_t elapsed = (start > stop) ? start - stop : stop - start;

{'fraction_non_alphanumeric': 0.1367709306800066, 'fraction_numerical': 0.005167390467813754, 'mean_word_length': 2.682591093117409, 'pattern_counts': {'":': 0, '<': 46, '<?xml version=': 0, '>': 46, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '6242848', 'n_tokens_mistral': 5442, 'n_tokens_neox': 4973, 'n_words': 1597}