<commit_before>/*
Copyright 2016 Fixstars Corporation

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

http ://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/

#include <cstdio>
#include <libsgm.h>
#include "winner_takes_all.hpp"
#include "utility.hpp"

namespace sgm {

namespace {

static constexpr unsigned int NUM_PATHS = 8u;

static constexpr unsigned int WARPS_PER_BLOCK = 8u;
static constexpr unsigned int BLOCK_SIZE = WARPS_PER_BLOCK * WARP_SIZE;


__device__ inline void update_top2(uint32_t& v0, uint32_t& v1, uint32_t x){
	const uint32_t y = max(x, v0);
	v0 = min(x, v0);
	v1 = min(y, v1);
}

struct Top2 {
	uint32_t values[2];

	__device__ void initialize(){
		values[0] = 0xffffffffu;
		values[1] = 0xffffffffu;
	}

	__device__ void push(uint32_t x){
		update_top2(values[0], values[1], x);
	}
};

template <unsigned int GROUP_SIZE, unsigned int STEP>
struct subgroup_merge_top2_impl {
	static __device__ Top2 call(Top2 x){
#if CUDA_VERSION >= 9000
		const uint32_t a = __shfl_xor_sync(0xffffffffu, x.values[0], STEP / 2, GROUP_SIZE);
		const uint32_t b = __shfl_xor_sync(0xffffffffu, x.values[1], STEP / 2, GROUP_SIZE);
#else
		const uint32_t a = __shfl_xor(x.values[0], STEP / 2, GROUP_SIZE);
		const uint32_t b = __shfl_xor(x.values[1], STEP / 2, GROUP_SIZE);
#endif
		x.push(a);
		x.push(b);
		return subgroup_merge_top2_impl<GROUP_SIZE, STEP / 2>::call(x);
	}
};

template <unsigned int GROUP_SIZE>
struct subgroup_merge_top2_impl<GROUP_SIZE, 1u> {
	static __device__ Top2 call(Top2 x){
		return x;
	}
};

template <unsigned int GROUP_SIZE>
__device__ inline Top2 subgroup_merge_top2(Top2 x){
	return subgroup_merge_top2_impl<GROUP_SIZE, GROUP_SIZE>::call(x);
}

__device__ inline uint32_t pack_cost_index(uint32_t cost, uint32_t index){
	union {
		uint32_t uint32;
		ushort2 uint16x2;
	} u;
	u.uint16x2.x = static_cast<uint16_t>(index);
	u.uint16x2.y = static_cast<uint16_t>(cost);
	return u.uint32;
}

__device__ uint32_t unpack_cost(uint32_t packed){
	return packed >> 16;
}

__device__ uint32_t unpack_index(uint32_t packed){
	return packed & 0xffffu;
}

using ComputeDisparity = uint16_t(*)(Top2, float, uint16_t*);

template <size_t MAX_DISPARITY>
__device__ inline uint16_t compute_disparity_normal(Top2 t2, float uniqueness, uint16_t* smem = nullptr)
{
	const float cost0 = static_cast<float>(unpack_cost(t2.values[0]));
	const float cost1 = static_cast<float>(unpack_cost(t2.values[1]));
	const uint16_t disp0 = static_cast<uint16_t>(unpack_index(t2.values[0]));
	const uint16_t disp1 = static_cast<uint16_t>(unpack_index(t2.values[1]));
	if(cost1 * uniqueness >= cost0){
		return disp0;
	}else if(abs(disp1 - disp0) <= 1){
		return disp0;
	}else{
		return 0;
	}
}

template <size_t MAX_DISPARITY>
__device__ inline uint16_t compute_disparity_subpixel(Top2 t2, float uniqueness, uint16_t* smem)
{
	uint16_t disp = compute_disparity_normal<MAX_DISPARITY>(t2, uniqueness, smem);
	disp <<= sgm::StereoSGM::SUBPIXEL_SHIFT;
	if (disp > 0 && disp < MAX_DISPARITY - 1) {
		const int numer = smem[disp - 1] - smem[disp + 1];
		const int denom = smem[disp - 1] - 2 * smem[disp] + smem[disp + 1];
		disp += (numer << sgm::StereoSGM::SUBPIXEL_SHIFT + denom) / (2 * denom);
	}
	return disp;
}


template <unsigned int MAX_DISPARITY, ComputeDisparity compute_disparity = compute_disparity_normal<MAX_DISPARITY>>
__global__ void winner_takes_all_kernel(
	output_type *left_dest,<commit_msg>Fix compute subpixel disparity<commit_after>/*
Copyright 2016 Fixstars Corporation

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

http ://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/

#include <cstdio>
#include <libsgm.h>
#include "winner_takes_all.hpp"
#include "utility.hpp"

namespace sgm {

namespace {

static constexpr unsigned int NUM_PATHS = 8u;

static constexpr unsigned int WARPS_PER_BLOCK = 8u;
static constexpr unsigned int BLOCK_SIZE = WARPS_PER_BLOCK * WARP_SIZE;


__device__ inline void update_top2(uint32_t& v0, uint32_t& v1, uint32_t x){
	const uint32_t y = max(x, v0);
	v0 = min(x, v0);
	v1 = min(y, v1);
}

struct Top2 {
	uint32_t values[2];

	__device__ void initialize(){
		values[0] = 0xffffffffu;
		values[1] = 0xffffffffu;
	}

	__device__ void push(uint32_t x){
		update_top2(values[0], values[1], x);
	}
};

template <unsigned int GROUP_SIZE, unsigned int STEP>
struct subgroup_merge_top2_impl {
	static __device__ Top2 call(Top2 x){
#if CUDA_VERSION >= 9000
		const uint32_t a = __shfl_xor_sync(0xffffffffu, x.values[0], STEP / 2, GROUP_SIZE);
		const uint32_t b = __shfl_xor_sync(0xffffffffu, x.values[1], STEP / 2, GROUP_SIZE);
#else
		const uint32_t a = __shfl_xor(x.values[0], STEP / 2, GROUP_SIZE);
		const uint32_t b = __shfl_xor(x.values[1], STEP / 2, GROUP_SIZE);
#endif
		x.push(a);
		x.push(b);
		return subgroup_merge_top2_impl<GROUP_SIZE, STEP / 2>::call(x);
	}
};

template <unsigned int GROUP_SIZE>
struct subgroup_merge_top2_impl<GROUP_SIZE, 1u> {
	static __device__ Top2 call(Top2 x){
		return x;
	}
};

template <unsigned int GROUP_SIZE>
__device__ inline Top2 subgroup_merge_top2(Top2 x){
	return subgroup_merge_top2_impl<GROUP_SIZE, GROUP_SIZE>::call(x);
}

__device__ inline uint32_t pack_cost_index(uint32_t cost, uint32_t index){
	union {
		uint32_t uint32;
		ushort2 uint16x2;
	} u;
	u.uint16x2.x = static_cast<uint16_t>(index);
	u.uint16x2.y = static_cast<uint16_t>(cost);
	return u.uint32;
}

__device__ uint32_t unpack_cost(uint32_t packed){
	return packed >> 16;
}

__device__ uint32_t unpack_index(uint32_t packed){
	return packed & 0xffffu;
}

using ComputeDisparity = uint16_t(*)(Top2, float, uint16_t*);

template <size_t MAX_DISPARITY>
__device__ inline uint16_t compute_disparity_normal(Top2 t2, float uniqueness, uint16_t* smem = nullptr)
{
	const float cost0 = static_cast<float>(unpack_cost(t2.values[0]));
	const float cost1 = static_cast<float>(unpack_cost(t2.values[1]));
	const uint16_t disp0 = static_cast<uint16_t>(unpack_index(t2.values[0]));
	const uint16_t disp1 = static_cast<uint16_t>(unpack_index(t2.values[1]));
	if(cost1 * uniqueness >= cost0){
		return disp0;
	}else if(abs(disp1 - disp0) <= 1){
		return disp0;
	}else{
		return 0;
	}
}

template <size_t MAX_DISPARITY>
__device__ inline uint16_t compute_disparity_subpixel(Top2 t2, float uniqueness, uint16_t* smem)
{
	const float cost0 = static_cast<float>(unpack_cost(t2.values[0]));
	const float cost1 = static_cast<float>(unpack_cost(t2.values[1]));
	const uint16_t disp0 = static_cast<uint16_t>(unpack_index(t2.values[0]));
	const uint16_t disp1 = static_cast<uint16_t>(unpack_index(t2.values[1]));
	if(cost1 * uniqueness >= cost0
		|| abs(disp1 - disp0) <= 1){
		uint16_t disp = disp0;
		disp <<= sgm::StereoSGM::SUBPIXEL_SHIFT;
		if (disp > 0 && disp < MAX_DISPARITY - 1) {
			const int numer = smem[disp - 1] - smem[disp + 1];
			const int denom = smem[disp - 1] - 2 * smem[disp] + smem[disp + 1];
			if (denom != 0 && -denom <= numer && numer <= denom) {
				disp += ((numer << sgm::StereoSGM::SUBPIXEL_SHIFT) + denom) / (2 * denom);
			}
		}
		return disp;
	}else{
		return 0;
	}
}


template <unsigned int MAX_DISPARITY, ComputeDisparity compute_disparity = compute_disparity_normal<MAX_DISPARITY>>
__global__ void winner_takes_all_kernel(
	output_type *left_dest,
{'fraction_non_alphanumeric': 0.1127530113001366, 'fraction_numerical': 0.044828014404569724, 'mean_word_length': 5.316862745098039, 'pattern_counts': {'":': 0, '<': 59, '<?xml version=': 0, '>': 55, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 8, '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': '6165903', 'n_tokens_mistral': 3323, 'n_tokens_neox': 3014, 'n_words': 782}