<commit_before>/******************************************************************************
 * Copyright (c) 2013, 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.
 *
 ******************************************************************************/

/******************************************************************************
 *
 * Code and text by <NAME>, NVIDIA Research
 * See http://nvlabs.github.io/moderngpu for repository and documentation.
 *
 ******************************************************************************/

#pragma once

#include "../mgpuhost.cuh"
#include "../device/launchbox.cuh"
#include "../device/loadstore.cuh"
#include "../device/ctaloadbalance.cuh"
#include "../kernels/search.cuh"

namespace mgpu {

////////////////////////////////////////////////////////////////////////////////
// KernelIntervalExpand

template<typename Tuning, typename IndicesIt, typename ValuesIt,
	typename OutputIt>
MGPU_LAUNCH_BOUNDS void KernelIntervalExpand(int destCount, 
	IndicesIt indices_global, ValuesIt values_global, int sourceCount, 
	const int* mp_global, OutputIt output_global) {

	typedef MGPU_LAUNCH_PARAMS Params;
	const int NT = Params::NT;
	const int VT = Params::VT;
	typedef typename std::iterator_traits<ValuesIt>::value_type T;

	union Shared {
		int indices[NT * (VT + 1)];
		T values[NT * VT];
	};
	__shared__ Shared shared;
	int tid = threadIdx.x;
	int block = blockIdx.x;

	// Compute the input and output intervals this CTA processes.
	int4 range = CTALoadBalance<NT, VT>(destCount, indices_global, sourceCount,
		block, tid, mp_global, shared.indices, true);
	
	// The interval indices are in the left part of shared memory (moveCount).
	// The scan of interval counts are in the right part (intervalCount).
	destCount = range.y - range.x;
	sourceCount = range.w - range.z;

	// Copy the source indices into register.
	int sources[VT];
	DeviceSharedToReg<NT, VT>(shared.indices, tid, sources);

	// Load the source fill values into shared memory. Each value is fetched
	// only once to reduce latency and L2 traffic.
	DeviceMemToMemLoop<NT>(sourceCount, values_global + range.z, tid,
		shared.values);
	
	// Gather the values from shared memory into register. This uses a shared
	// memory broadcast - one instance of a value serves all the threads that
	// comprise its fill operation.
	T values[VT];
	DeviceGather<NT, VT>(destCount, shared.values - range.z, sources, tid,
		values, false);

	// Store the values to global memory.
	DeviceRegToGlobal<NT, VT>(destCount, values, tid, output_global + range.x);
}

////////////////////////////////////////////////////////////////////////////////
// IntervalExpand

template<typename IndicesIt, typename ValuesIt, typename OutputIt>
MGPU_HOST void IntervalExpand(int moveCount, IndicesIt indices_global, 
	ValuesIt values_global, int intervalCount, OutputIt output_global,
	CudaContext& context) {

	const int NT = 128;
	const int VT = 7;
	typedef LaunchBoxVT<NT, VT> Tuning;<commit_msg>Fix a compiler warning in ModernGPU<commit_after>/******************************************************************************
 * Copyright (c) 2013, 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.
 *
 ******************************************************************************/

/******************************************************************************
 *
 * Code and text by <NAME>, NVIDIA Research
 * See http://nvlabs.github.io/moderngpu for repository and documentation.
 *
 ******************************************************************************/

#pragma once

#include "../mgpuhost.cuh"
#include "../device/launchbox.cuh"
#include "../device/loadstore.cuh"
#include "../device/ctaloadbalance.cuh"
#include "../kernels/search.cuh"

namespace mgpu {

////////////////////////////////////////////////////////////////////////////////
// KernelIntervalExpand

template<typename Tuning, typename IndicesIt, typename ValuesIt,
	typename OutputIt>
MGPU_LAUNCH_BOUNDS void KernelIntervalExpand(int destCount, 
	IndicesIt indices_global, ValuesIt values_global, int sourceCount, 
	const int* mp_global, OutputIt output_global) {

	typedef MGPU_LAUNCH_PARAMS Params;
	const int NT = Params::NT;
	const int VT = Params::VT;
	typedef typename std::iterator_traits<ValuesIt>::value_type T;

	union Shared {
		int indices[NT * (VT + 1)];
		T values[NT * VT];
	};
	__shared__ Shared shared;
	int tid = threadIdx.x;
	int block = blockIdx.x;

	// Compute the input and output intervals this CTA processes.
	int4 range = CTALoadBalance<NT, VT>(destCount, indices_global, sourceCount,
		block, tid, mp_global, shared.indices, true);
	
	// The interval indices are in the left part of shared memory (moveCount).
	// The scan of interval counts are in the right part (intervalCount).
	destCount = range.y - range.x;
	sourceCount = range.w - range.z;

	// Copy the source indices into register.
	int sources[VT];
	DeviceSharedToReg<NT, VT>(shared.indices, tid, sources);

	// Load the source fill values into shared memory. Each value is fetched
	// only once to reduce latency and L2 traffic.
	DeviceMemToMemLoop<NT>(sourceCount, values_global + range.z, tid,
		shared.values);
	
	// Gather the values from shared memory into register. This uses a shared
	// memory broadcast - one instance of a value serves all the threads that
	// comprise its fill operation.
	T values[VT];
	DeviceGather<NT, VT>(destCount, shared.values - range.z, sources, tid,
		values, false);

	// Store the values to global memory.
	DeviceRegToGlobal<NT, VT>(destCount, values, tid, output_global + range.x);
}

////////////////////////////////////////////////////////////////////////////////
// IntervalExpand

template<typename IndicesIt, typename ValuesIt, typename OutputIt>
MGPU_HOST void IntervalExpand(int moveCount, IndicesIt indices_global, 
	ValuesIt values_global, int intervalCount, OutputIt output_global,
	CudaContext& context) {

	//const int NT = 128;
	//const int VT = 7;
	typedef LaunchBoxVT<128, 7> Tuning;
{'fraction_non_alphanumeric': 0.17543859649122806, 'fraction_numerical': 0.0028869642460581835, 'mean_word_length': 3.7656084656084654, 'pattern_counts': {'":': 0, '<': 23, '<?xml version=': 0, '>': 23, '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': '2259379', 'n_tokens_mistral': 2480, 'n_tokens_neox': 2233, 'n_words': 948}