hyzboy/TexConv
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
TexConv/CMP_Core/shaders/bcn_common_api.h

1423 lines
37 KiB
C++
Raw Permalink Blame History

//===============================================================================
// Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files(the "Software"), to deal
// in the Software without restriction, including without limitation the rights to
// use, copy, modify, merge, publish, distribute, sublicense, and / or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions :
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.
//
//===============================================================================
#ifndef BCN_COMMON_API_H_
#define BCN_COMMON_API_H_
//===================================================================
// NOTE: Do not use these API in production code, subject to changes
//===================================================================
#ifndef ASPM_GPU
#pragma warning(disable : 4244)
#pragma warning(disable : 4201)
#endif
#if defined(_M_X64) || defined(_M_IX86) || defined(x86_64) || defined(i386)
#include <xmmintrin.h>
#include <smmintrin.h>
#elif defined(arm) || defined(_M_ARM) || defined(aarch64)
#include <arm_neon.h>
#endif
#include "common_def.h"
#define CMP_MAX_16BITFLOAT 65504.0f
#define CMP_FLT_MAX 3.402823466e+38F
#define BC1ConstructColour(r, g, b) (((r) << 11) | ((g) << 5) | (b))
#ifdef ASPM_HLSL
#define fabs(x) abs(x)
#endif
CMP_STATIC CGU_FLOAT cmp_fabs(CMP_IN CGU_FLOAT x)
{
return fabs(x);
}
CMP_STATIC CGU_FLOAT cmp_linearToSrgbf(CMP_IN CGU_FLOAT Color)
{
if (Color <= 0.0f)
return (0.0f);
if (Color >= 1.0f)
return (1.0f);
// standard : 0.0031308f
if (Color <= 0.00313066844250063)
return (Color * 12.92f);
return (pow(fabs(Color), 1.0f / 2.4f) * 1.055f - 0.055f);
}
CMP_STATIC CGU_Vec3f cmp_linearToSrgb(CMP_IN CGU_Vec3f Color)
{
Color.x = cmp_linearToSrgbf(Color.x);
Color.y = cmp_linearToSrgbf(Color.y);
Color.z = cmp_linearToSrgbf(Color.z);
return Color;
}
CMP_STATIC CGU_FLOAT cmp_srgbToLinearf(CMP_IN CGU_FLOAT Color)
{
if (Color <= 0.0f)
return (0.0f);
if (Color >= 1.0f)
return (1.0f);
// standard 0.04045f
if (Color <= 0.0404482362771082)
return (Color / 12.92f);
return pow((Color + 0.055f) / 1.055f, 2.4f);
}
CMP_STATIC CGU_Vec3f cmp_srgbToLinear(CMP_IN CGU_Vec3f Color)
{
Color.x = cmp_srgbToLinearf(Color.x);
Color.y = cmp_srgbToLinearf(Color.y);
Color.z = cmp_srgbToLinearf(Color.z);
return Color;
}
CMP_STATIC CGU_Vec3f cmp_565ToLinear(CMP_IN CGU_UINT32 n565)
{
CGU_UINT32 r0;
CGU_UINT32 g0;
CGU_UINT32 b0;
r0 = ((n565 & 0xf800) >> 8);
g0 = ((n565 & 0x07e0) >> 3);
b0 = ((n565 & 0x001f) << 3);
// Apply the lower bit replication to give full dynamic range (5,6,5)
r0 += (r0 >> 5);
g0 += (g0 >> 6);
b0 += (b0 >> 5);
CGU_Vec3f LinearColor;
LinearColor.x = (CGU_FLOAT)r0;
LinearColor.y = (CGU_FLOAT)g0;
LinearColor.z = (CGU_FLOAT)b0;
return LinearColor;
}
CMP_STATIC CGU_UINT32 cmp_get2Bit32(CMP_IN CGU_UINT32 value, CMP_IN CGU_UINT32 indexPos)
{
return (value >> (indexPos * 2)) & 0x3;
}
CMP_STATIC CGU_UINT32 cmp_set2Bit32(CMP_IN CGU_UINT32 value, CMP_IN CGU_UINT32 indexPos)
{
return ((value & 0x3) << (indexPos * 2));
}
CMP_STATIC CGU_UINT32 cmp_constructColor(CMP_IN CGU_UINT32 R, CMP_IN CGU_UINT32 G, CMP_IN CGU_UINT32 B)
{
return (((R & 0x000000F8) << 8) | ((G & 0x000000FC) << 3) | ((B & 0x000000F8) >> 3));
}
CMP_STATIC CGU_FLOAT cmp_clampf(CMP_IN CGU_FLOAT v, CMP_IN CGU_FLOAT a, CMP_IN CGU_FLOAT b)
{
if (v < a)
return a;
else if (v > b)
return b;
return v;
}
CMP_STATIC CGU_Vec3f cmp_clampVec3f(CMP_IN CGU_Vec3f value, CMP_IN CGU_FLOAT minValue, CMP_IN CGU_FLOAT maxValue)
{
#ifdef ASPM_GPU
return clamp(value, minValue, maxValue);
#else
CGU_Vec3f revalue;
revalue.x = cmp_clampf(value.x, minValue, maxValue);
revalue.y = cmp_clampf(value.y, minValue, maxValue);
revalue.z = cmp_clampf(value.z, minValue, maxValue);
return revalue;
#endif
}
CMP_STATIC CGU_Vec3f cmp_saturate(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_HLSL
return saturate(value);
#else
return cmp_clampVec3f(value, 0.0f, 1.0f);
#endif
}
static CGU_Vec3f cmp_powVec3f(CGU_Vec3f color, CGU_FLOAT ex)
{
#ifdef ASPM_GPU
return pow(color, ex);
#else
CGU_Vec3f ColorSrgbPower;
ColorSrgbPower.x = pow(color.x, ex);
ColorSrgbPower.y = pow(color.y, ex);
ColorSrgbPower.z = pow(color.z, ex);
return ColorSrgbPower;
#endif
}
CMP_STATIC CGU_Vec3f cmp_minVec3f(CMP_IN CGU_Vec3f a, CMP_IN CGU_Vec3f b)
{
#ifdef ASPM_HLSL
return min(a, b);
#endif
CGU_Vec3f res;
if (a.x < b.x)
res.x = a.x;
else
res.x = b.x;
if (a.y < b.y)
res.y = a.y;
else
res.y = b.y;
if (a.z < b.z)
res.z = a.z;
else
res.z = b.z;
return res;
}
CMP_STATIC CGU_Vec3f cmp_maxVec3f(CMP_IN CGU_Vec3f a, CMP_IN CGU_Vec3f b)
{
#ifdef ASPM_HLSL
return max(a, b);
#endif
CGU_Vec3f res;
if (a.x > b.x)
res.x = a.x;
else
res.x = b.x;
if (a.y > b.y)
res.y = a.y;
else
res.y = b.y;
if (a.z > b.z)
res.z = a.z;
else
res.z = b.z;
return res;
}
inline CGU_Vec3f cmp_min3f(CMP_IN CGU_Vec3f value1, CMP_IN CGU_Vec3f value2)
{
#ifdef ASPM_GPU
return min(value1, value2);
#else
CGU_Vec3f res;
res.x = CMP_MIN(value1.x, value2.x);
res.y = CMP_MIN(value1.y, value2.y);
res.z = CMP_MIN(value1.z, value2.z);
return res;
#endif
}
inline CGU_Vec3f cmp_max3f(CMP_IN CGU_Vec3f value1, CMP_IN CGU_Vec3f value2)
{
#ifdef ASPM_GPU
return max(value1, value2);
#else
CGU_Vec3f res;
res.x = CMP_MAX(value1.x, value2.x);
res.y = CMP_MAX(value1.y, value2.y);
res.z = CMP_MAX(value1.z, value2.z);
return res;
#endif
}
CMP_STATIC CGU_FLOAT cmp_minf(CMP_IN CGU_FLOAT a, CMP_IN CGU_FLOAT b)
{
return a < b ? a : b;
}
CMP_STATIC CGU_FLOAT cmp_maxf(CMP_IN CGU_FLOAT a, CMP_IN CGU_FLOAT b)
{
return a > b ? a : b;
}
CMP_STATIC CGU_FLOAT cmp_floor(CMP_IN CGU_FLOAT value)
{
#if defined(_M_X64) || defined(_M_IX86) || defined(x86_64) || defined(i386)
__m128 vectorValue = _mm_set_ss(value);
return _mm_cvtss_f32(_mm_floor_ss(vectorValue, vectorValue));
#elif defined(arm) || defined(_M_ARM) || defined(aarch64)
return vget_lane_f32(vrndm_f32(vdup_n_f32(value)), 0);
#else
return floor(value);
#endif
}
CMP_STATIC CGU_Vec3f cmp_floorVec3f(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_GPU
return floor(value);
#else
CGU_Vec3f revalue;
revalue.x = floor(value.x);
revalue.y = floor(value.y);
revalue.z = floor(value.z);
return revalue;
#endif
}
#ifndef ASPM_OPENCL
//=======================================================
// COMMON GPU & CPU API
//=======================================================
//======================
// implicit vector cast
//======================
CMP_STATIC CGU_Vec4i cmp_castimp(CGU_Vec4ui v1)
{
#ifdef ASPM_HLSL
return (v1);
#else
return (v1.x, v1.y, v1.z, v1.w);
#endif
}
CMP_STATIC CGU_Vec3i cmp_castimp(CGU_Vec3ui v1)
{
#ifdef ASPM_HLSL
return (v1);
#else
return (v1.x, v1.y, v1.z);
#endif
}
//======================
// Min / Max
//======================
CMP_STATIC CGU_UINT8 cmp_min8(CMP_IN CGU_UINT8 a, CMP_IN CGU_UINT8 b)
{
return a < b ? a : b;
}
CMP_STATIC CGU_UINT8 cmp_max8(CMP_IN CGU_UINT8 a, CMP_IN CGU_UINT8 b)
{
return a > b ? a : b;
}
CMP_STATIC CGU_UINT32 cmp_mini(CMP_IN CGU_UINT32 a, CMP_IN CGU_UINT32 b)
{
return (a < b) ? a : b;
}
CMP_STATIC CGU_UINT32 cmp_maxi(CMP_IN CGU_UINT32 a, CMP_IN CGU_UINT32 b)
{
return (a > b) ? a : b;
}
CMP_STATIC CGU_FLOAT cmp_max3(CMP_IN CGU_FLOAT i, CMP_IN CGU_FLOAT j, CMP_IN CGU_FLOAT k)
{
#ifdef ASPM_GLSL
return max3(i, j, k);
#else
CGU_FLOAT max = i;
if (max < j)
max = j;
if (max < k)
max = k;
return (max);
#endif
}
CMP_STATIC CGU_Vec4ui cmp_minVec4ui(CMP_IN CGU_Vec4ui a, CMP_IN CGU_Vec4ui b)
{
//#ifdef ASPM_HLSL
// return min(a, b);
//#endif
//#ifndef ASPM_GPU
CGU_Vec4ui res;
if (a.x < b.x)
res.x = a.x;
else
res.x = b.x;
if (a.y < b.y)
res.y = a.y;
else
res.y = b.y;
if (a.z < b.z)
res.z = a.z;
else
res.z = b.z;
if (a.w < b.w)
res.w = a.w;
else
res.w = b.w;
return res;
//#endif
}
CMP_STATIC CGU_Vec4ui cmp_maxVec4ui(CMP_IN CGU_Vec4ui a, CMP_IN CGU_Vec4ui b)
{
//#ifdef ASPM_HLSL
// return max(a, b);
//#endif
//#ifndef ASPM_GPU
CGU_Vec4ui res;
if (a.x > b.x)
res.x = a.x;
else
res.x = b.x;
if (a.y > b.y)
res.y = a.y;
else
res.y = b.y;
if (a.z > b.z)
res.z = a.z;
else
res.z = b.z;
if (a.w > b.w)
res.w = a.w;
else
res.w = b.w;
return res;
//#endif
}
//======================
// Clamps
//======================
CMP_STATIC CGU_UINT32 cmp_clampui32(CMP_IN CGU_UINT32 v, CMP_IN CGU_UINT32 a, CMP_IN CGU_UINT32 b)
{
if (v < a)
return a;
else if (v > b)
return b;
return v;
}
// Test Ref:https://en.wikipedia.org/wiki/Half-precision_floating-point_format
// Half (in Hex) Float Comment
// ---------------------------------------------------------------------------
// 0001 (approx) = 0.000000059604645 smallest positive subnormal number
// 03ff (approx) = 0.000060975552 largest subnormal number
// 0400 (approx) = 0.00006103515625 smallest positive normal number
// 7bff (approx) = 65504 largest normal number
// 3bff (approx) = 0.99951172 largest number less than one
// 3c00 (approx) = 1.00097656 smallest number larger than one
// 3555 = 0.33325195 the rounding of 1/3 to nearest
// c000 = ?2
// 8000 = -0
// 0000 = 0
// 7c00 = infinity
// fc00 = infinity
// Half Float Math
CMP_STATIC CGU_FLOAT HalfToFloat(CGU_UINT32 h)
{
#if defined(ASPM_GPU)
CGU_FLOAT f = min16float((float)(h));
return f;
#else
union FP32
{
CGU_UINT32 u;
CGU_FLOAT f;
};
const FP32 magic = {(254 - 15) << 23};
const FP32 was_infnan = {(127 + 16) << 23};
FP32 o;
o.u = (h & 0x7fff) << 13; // exponent/mantissa bits
o.f *= magic.f; // exponent adjust
if (o.f >= was_infnan.f) // check Inf/NaN
o.u |= 255 << 23;
o.u |= (h & 0x8000) << 16; // sign bit
return o.f;
#endif
}
// From BC6HEcode.hlsl
CMP_STATIC CGU_FLOAT cmp_half2float1(CGU_UINT32 Value)
{
CGU_UINT32 Mantissa = (CGU_UINT32)(Value & 0x03FF);
CGU_UINT32 Exponent;
if ((Value & 0x7C00) != 0) // The value is normalized
{
Exponent = (CGU_UINT32)((Value >> 10) & 0x1F);
}
else if (Mantissa != 0) // The value is denormalized
{
// Normalize the value in the resulting float
Exponent = 1;
do
{
Exponent--;
Mantissa <<= 1;
} while ((Mantissa & 0x0400) == 0);
Mantissa &= 0x03FF;
}
else // The value is zero
{
Exponent = (CGU_UINT32)(-112);
}
CGU_UINT32 Result = ((Value & 0x8000) << 16) | // Sign
((Exponent + 112) << 23) | // Exponent
(Mantissa << 13); // Mantissa
return CGU_FLOAT(Result);
}
CMP_STATIC CGU_Vec3f cmp_half2floatVec3(CGU_Vec3ui color_h)
{
//uint3 sign = color_h & 0x8000;
//uint3 expo = color_h & 0x7C00;
//uint3 base = color_h & 0x03FF;
//return ( expo == 0 ) ? asfloat( ( sign << 16 ) | asuint( float3(base) / 16777216 ) ) //16777216 = 2^24
// : asfloat( ( sign << 16 ) | ( ( ( expo + 0x1C000 ) | base ) << 13 ) ); //0x1C000 = 0x1FC00 - 0x3C00
return CGU_Vec3f(cmp_half2float1(color_h.x), cmp_half2float1(color_h.y), cmp_half2float1(color_h.z));
}
CMP_STATIC CGU_UINT16 FloatToHalf(CGU_FLOAT value)
{
#if defined(ASPM_GPU)
return 0;
#else
union FP32
{
CGU_UINT16 u;
float f;
struct
{
CGU_UINT32 Mantissa : 23;
CGU_UINT32 Exponent : 8;
CGU_UINT32 Sign : 1;
};
};
union FP16
{
CGU_UINT16 u;
struct
{
CGU_UINT32 Mantissa : 10;
CGU_UINT32 Exponent : 5;
CGU_UINT32 Sign : 1;
};
};
FP16 o = {0};
FP32 f;
f.f = value;
// Based on ISPC reference code (with minor modifications)
if (f.Exponent == 0) // Signed zero/denormal (which will underflow)
o.Exponent = 0;
else if (f.Exponent == 255) // Inf or NaN (all exponent bits set)
{
o.Exponent = 31;
o.Mantissa = f.Mantissa ? 0x200 : 0; // NaN->qNaN and Inf->Inf
}
else // Normalized number
{
// Exponent unbias the single, then bias the halfp
int newexp = f.Exponent - 127 + 15;
if (newexp >= 31) // Overflow, return signed infinity
o.Exponent = 31;
else if (newexp <= 0) // Underflow
{
if ((14 - newexp) <= 24) // Mantissa might be non-zero
{
CGU_UINT32 mant = f.Mantissa | 0x800000; // Hidden 1 bit
o.Mantissa = mant >> (14 - newexp);
if ((mant >> (13 - newexp)) & 1) // Check for rounding
o.u++; // Round, might overflow into exp bit, but this is OK
}
}
else
{
o.Exponent = newexp;
o.Mantissa = f.Mantissa >> 13;
if (f.Mantissa & 0x1000) // Check for rounding
o.u++; // Round, might overflow to inf, this is OK
}
}
o.Sign = f.Sign;
return o.u;
#endif
}
CMP_STATIC CGU_UINT32 cmp_float2halfui(CGU_FLOAT f)
{
CGU_UINT32 Result;
CGU_UINT32 IValue = CGU_UINT32(f);
CGU_UINT32 Sign = (IValue & 0x80000000U) >> 16U;
IValue = IValue & 0x7FFFFFFFU;
if (IValue > 0x47FFEFFFU)
{
// The number is too large to be represented as a half. Saturate to infinity.
Result = 0x7FFFU;
}
else
{
if (IValue < 0x38800000U)
{
// The number is too small to be represented as a normalized half.
// Convert it to a denormalized value.
CGU_UINT32 Shift = 113U - (IValue >> 23U);
IValue = (0x800000U | (IValue & 0x7FFFFFU)) >> Shift;
}
else
{
// Rebias the exponent to represent the value as a normalized half.
IValue += 0xC8000000U;
}
Result = ((IValue + 0x0FFFU + ((IValue >> 13U) & 1U)) >> 13U) & 0x7FFFU;
}
return (Result | Sign);
}
CMP_STATIC CGU_Vec3ui cmp_float2half(CGU_Vec3f endPoint_f)
{
return CGU_Vec3ui(cmp_float2halfui(endPoint_f.x), cmp_float2halfui(endPoint_f.y), cmp_float2halfui(endPoint_f.z));
}
CMP_STATIC CGU_UINT32 cmp_float2half1(CGU_FLOAT f)
{
CGU_UINT32 Result;
CGU_UINT32 IValue = CGU_UINT32(f); //asuint(f);
CGU_UINT32 Sign = (IValue & 0x80000000U) >> 16U;
IValue = IValue & 0x7FFFFFFFU;
if (IValue > 0x47FFEFFFU)
{
// The number is too large to be represented as a half. Saturate to infinity.
Result = 0x7FFFU;
}
else
{
if (IValue < 0x38800000U)
{
// The number is too small to be represented as a normalized half.
// Convert it to a denormalized value.
CGU_UINT32 Shift = 113U - (IValue >> 23U);
IValue = (0x800000U | (IValue & 0x7FFFFFU)) >> Shift;
}
else
{
// Rebias the exponent to represent the value as a normalized half.
IValue += 0xC8000000U;
}
Result = ((IValue + 0x0FFFU + ((IValue >> 13U) & 1U)) >> 13U) & 0x7FFFU;
}
return (Result | Sign);
}
CMP_STATIC CGU_Vec3ui cmp_float2halfVec3(CGU_Vec3f endPoint_f)
{
return CGU_Vec3ui(cmp_float2half1(endPoint_f.x), cmp_float2half1(endPoint_f.y), cmp_float2half1(endPoint_f.z));
}
CMP_STATIC CGU_FLOAT cmp_f32tof16(CMP_IN CGU_FLOAT value)
{
#ifdef ASPM_GLSL
return packHalf2x16(CGU_Vec2f(value.x, 0.0));
#endif
#ifdef ASPM_HLSL
return f32tof16(value);
#endif
#ifndef ASPM_GPU
return FloatToHalf(value);
#endif
}
CMP_STATIC CGU_Vec3f cmp_f32tof16(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_GLSL
return CGU_Vec3f(packHalf2x16(CGU_Vec2f(value.x, 0.0)), packHalf2x16(CGU_Vec2f(value.y, 0.0)), packHalf2x16(CGU_Vec2f(value.z, 0.0)));
#endif
#ifdef ASPM_HLSL
return f32tof16(value);
#endif
#ifndef ASPM_GPU
CGU_Vec3f res;
res.x = FloatToHalf(value.x);
res.y = FloatToHalf(value.y);
res.z = FloatToHalf(value.z);
return res;
#endif
}
CMP_STATIC CGU_FLOAT cmp_f16tof32(CGU_UINT32 value)
{
#ifdef ASPM_GLSL
return unpackHalf2x16(value).x;
#endif
#ifdef ASPM_HLSL
return f16tof32(value);
#endif
#ifndef ASPM_GPU
return HalfToFloat(value);
#endif
}
CMP_STATIC CGU_Vec3f cmp_f16tof32(CGU_Vec3ui value)
{
#ifdef ASPM_GLSL
return CGU_Vec3f(unpackHalf2x16(value.x).x, unpackHalf2x16(value.y).x, unpackHalf2x16(value.z).x);
#endif
#ifdef ASPM_HLSL
return f16tof32(value);
#endif
#ifndef ASPM_GPU
CGU_Vec3f res;
res.x = HalfToFloat(value.x);
res.y = HalfToFloat(value.y);
res.z = HalfToFloat(value.z);
return res;
#endif
}
CMP_STATIC CGU_Vec3f cmp_f16tof32(CGU_Vec3f value)
{
#ifdef ASPM_GLSL
return CGU_Vec3f(unpackHalf2x16(value.x).x, unpackHalf2x16(value.y).x, unpackHalf2x16(value.z).x);
#endif
#ifdef ASPM_HLSL
return f16tof32(value);
#endif
#ifndef ASPM_GPU
CGU_Vec3f res;
res.x = HalfToFloat((CGU_UINT32)value.x);
res.y = HalfToFloat((CGU_UINT32)value.y);
res.z = HalfToFloat((CGU_UINT32)value.z);
return res;
#endif
}
CMP_STATIC void cmp_swap(CMP_INOUT CGU_Vec3f CMP_REFINOUT a, CMP_INOUT CGU_Vec3f CMP_REFINOUT b)
{
CGU_Vec3f tmp = a;
a = b;
b = tmp;
}
CMP_STATIC void cmp_swap(CMP_INOUT CGU_FLOAT CMP_REFINOUT a, CMP_INOUT CGU_FLOAT CMP_REFINOUT b)
{
CGU_FLOAT tmp = a;
a = b;
b = tmp;
}
CMP_STATIC void cmp_swap(CMP_INOUT CGU_Vec3i CMP_REFINOUT lhs, CMP_INOUT CGU_Vec3i CMP_REFINOUT rhs) // valided with msc code
{
CGU_Vec3i tmp = lhs;
lhs = rhs;
rhs = tmp;
}
CMP_STATIC CGU_INT cmp_dotVec2i(CMP_IN CGU_Vec2i value1, CMP_IN CGU_Vec2i value2)
{
#ifdef ASPM_GPU
return dot(value1, value2);
#else
return (value1.x * value2.x) + (value1.y * value2.y);
#endif
}
CMP_STATIC CGU_FLOAT cmp_dotVec3f(CMP_IN CGU_Vec3f value1, CMP_IN CGU_Vec3f value2)
{
#ifdef ASPM_GPU
return dot(value1, value2);
#else
return (value1.x * value2.x) + (value1.y * value2.y) + (value1.z * value2.z);
#endif
}
CMP_STATIC CGU_UINT32 cmp_dotVec3ui(CMP_IN CGU_Vec3ui value1, CMP_IN CGU_Vec3ui value2)
{
#ifdef ASPM_GPU
return dot(value1, value2);
#else
return (value1.x * value2.x) + (value1.y * value2.y) + (value1.z * value2.z);
#endif
}
CMP_STATIC CGU_UINT32 cmp_dotVec4i(CMP_IN CGU_Vec4i value1, CMP_IN CGU_Vec4i value2)
{
#ifdef ASPM_GPU
return dot(value1, value2);
#else
return (value1.x * value2.x) + (value1.y * value2.y) + (value1.z * value2.z) + (value1.w * value2.w);
#endif
}
CMP_STATIC CGU_UINT32 cmp_dotVec4ui(CMP_IN CGU_Vec4ui value1, CMP_IN CGU_Vec4ui value2)
{
#ifdef ASPM_GPU
return dot(value1, value2);
#else
return (value1.x * value2.x) + (value1.y * value2.y) + (value1.z * value2.z) + (value1.w * value2.w);
#endif
}
CMP_STATIC CGU_Vec3f cmp_clampVec3fi(CMP_IN CGU_Vec3f value, CMP_IN CGU_INT minValue, CMP_IN CGU_INT maxValue)
{
#ifdef ASPM_GPU
return clamp(value, minValue, maxValue);
#else
CGU_Vec3f revalue;
revalue.x = cmp_clampf(value.x, (CGU_FLOAT)minValue, (CGU_FLOAT)maxValue);
revalue.y = cmp_clampf(value.y, (CGU_FLOAT)minValue, (CGU_FLOAT)maxValue);
revalue.z = cmp_clampf(value.z, (CGU_FLOAT)minValue, (CGU_FLOAT)maxValue);
return revalue;
#endif
}
CMP_STATIC CGU_Vec4ui cmp_clampVec4ui(CMP_IN CGU_Vec4ui value, CMP_IN CGU_UINT32 minValue, CMP_IN CGU_UINT32 maxValue)
{
#ifdef ASPM_GPU
return clamp(value, minValue, maxValue);
#else
CGU_Vec4ui revalue;
revalue.x = cmp_clampui32(value.x, minValue, maxValue);
revalue.y = cmp_clampui32(value.y, minValue, maxValue);
revalue.z = cmp_clampui32(value.z, minValue, maxValue);
revalue.w = cmp_clampui32(value.w, minValue, maxValue);
return revalue;
#endif
}
CMP_STATIC CGU_Vec4f cmp_clampVec4f(CMP_IN CGU_Vec4f value, CMP_IN CGU_FLOAT minValue, CMP_IN CGU_FLOAT maxValue)
{
#ifdef ASPM_GPU
return clamp(value, minValue, maxValue);
#else
CGU_Vec4f revalue;
revalue.x = cmp_clampf(value.x, minValue, maxValue);
revalue.y = cmp_clampf(value.y, minValue, maxValue);
revalue.z = cmp_clampf(value.z, minValue, maxValue);
revalue.w = cmp_clampf(value.w, minValue, maxValue);
return revalue;
#endif
}
CMP_STATIC CGU_Vec3f cmp_clamp3Vec3f(CMP_IN CGU_Vec3f value, CMP_IN CGU_Vec3f minValue, CMP_IN CGU_Vec3f maxValue)
{
#ifdef ASPM_GPU
return clamp(value, minValue, maxValue);
#else
CGU_Vec3f revalue;
revalue.x = cmp_clampf(value.x, minValue.x, maxValue.x);
revalue.y = cmp_clampf(value.y, minValue.y, maxValue.y);
revalue.z = cmp_clampf(value.z, minValue.z, maxValue.z);
return revalue;
#endif
}
CMP_STATIC CGU_Vec3f cmp_exp2(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_GPU
return exp2(value);
#else
CGU_Vec3f revalue;
revalue.x = exp2(value.x);
revalue.y = exp2(value.y);
revalue.z = exp2(value.z);
return revalue;
#endif
}
CMP_STATIC CGU_Vec3f cmp_roundVec3f(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_HLSL
return round(value);
#endif
#ifndef ASPM_HLSL
CGU_Vec3f res;
res.x = round(value.x);
res.y = round(value.y);
res.z = round(value.z);
return res;
#endif
}
CMP_STATIC CGU_Vec3f cmp_log2Vec3f(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_GPU
return log2(value);
#else
CGU_Vec3f res;
res.x = log2(value.x);
res.y = log2(value.y);
res.z = log2(value.z);
return res;
#endif
}
// used in BC1 LowQuality code
CMP_STATIC CGU_FLOAT cmp_saturate(CMP_IN CGU_FLOAT value)
{
#ifdef ASPM_HLSL
return saturate(value);
#else
return cmp_clampf(value, 0.0f, 1.0f);
#endif
}
CMP_STATIC CGU_FLOAT cmp_rcp(CMP_IN CGU_FLOAT det)
{
#ifdef ASPM_HLSL
return rcp(det);
#else
if (det > 0.0f)
return (1 / det);
else
return 0.0f;
#endif
}
CMP_STATIC CGU_UINT32 cmp_Get4BitIndexPos(CMP_IN CGU_FLOAT indexPos, CMP_IN CGU_FLOAT endPoint0Pos, CMP_IN CGU_FLOAT endPoint1Pos)
{
CGU_FLOAT r = (indexPos - endPoint0Pos) / (endPoint1Pos - endPoint0Pos);
return cmp_clampui32(CGU_UINT32(r * 14.93333f + 0.03333f + 0.5f), 0, 15);
}
// Calculate Mean Square Least Error (MSLE) for 2 Vectors
CMP_STATIC CGU_FLOAT cmp_CalcMSLE(CMP_IN CGU_Vec3f a, CMP_IN CGU_Vec3f b)
{
CGU_Vec3f err = cmp_log2Vec3f((b + 1.0f) / (a + 1.0f));
err = err * err;
return err.x + err.y + err.z;
}
// Compute Endpoints (min/max) bounding box
CMP_STATIC void cmp_GetTexelMinMax(CMP_IN CGU_Vec3f texels[16], CMP_INOUT CGU_Vec3f CMP_REFINOUT blockMin, CMP_INOUT CGU_Vec3f CMP_REFINOUT blockMax)
{
blockMin = texels[0];
blockMax = texels[0];
for (CGU_UINT32 i = 1u; i < 16u; ++i)
{
blockMin = cmp_minVec3f(blockMin, texels[i]);
blockMax = cmp_maxVec3f(blockMax, texels[i]);
}
}
// Refine Endpoints (min/max) by insetting bounding box in log2 RGB space
CMP_STATIC void cmp_RefineMinMaxAsLog2(CMP_IN CGU_Vec3f texels[16], CMP_INOUT CGU_Vec3f CMP_REFINOUT blockMin, CMP_INOUT CGU_Vec3f CMP_REFINOUT blockMax)
{
CGU_Vec3f refinedBlockMin = blockMax;
CGU_Vec3f refinedBlockMax = blockMin;
for (CGU_UINT32 i = 0u; i < 16u; ++i)
{
refinedBlockMin = cmp_minVec3f(refinedBlockMin, texels[i] == blockMin ? refinedBlockMin : texels[i]);
refinedBlockMax = cmp_maxVec3f(refinedBlockMax, texels[i] == blockMax ? refinedBlockMax : texels[i]);
}
CGU_Vec3f logBlockMax = cmp_log2Vec3f(blockMax + 1.0f);
CGU_Vec3f logBlockMin = cmp_log2Vec3f(blockMin + 1.0f);
CGU_Vec3f logRefinedBlockMax = cmp_log2Vec3f(refinedBlockMax + 1.0f);
CGU_Vec3f logRefinedBlockMin = cmp_log2Vec3f(refinedBlockMin + 1.0f);
CGU_Vec3f logBlockMaxExt = (logBlockMax - logBlockMin) * (1.0f / 32.0f);
logBlockMin += cmp_minVec3f(logRefinedBlockMin - logBlockMin, logBlockMaxExt);
logBlockMax -= cmp_minVec3f(logBlockMax - logRefinedBlockMax, logBlockMaxExt);
blockMin = cmp_exp2(logBlockMin) - 1.0f;
blockMax = cmp_exp2(logBlockMax) - 1.0f;
}
// Refine Endpoints (min/max) by Least Squares Optimization
CMP_STATIC void cmp_RefineMinMaxAs16BitLeastSquares(CMP_IN CGU_Vec3f texels[16],
CMP_INOUT CGU_Vec3f CMP_REFINOUT blockMin,
CMP_INOUT CGU_Vec3f CMP_REFINOUT blockMax)
{
CGU_Vec3f blockDir = blockMax - blockMin;
blockDir = blockDir / (blockDir.x + blockDir.y + blockDir.z);
CGU_FLOAT endPoint0Pos = cmp_f32tof16(cmp_dotVec3f(blockMin, blockDir));
CGU_FLOAT endPoint1Pos = cmp_f32tof16(cmp_dotVec3f(blockMax, blockDir));
CGU_Vec3f alphaTexelSum = 0.0f;
CGU_Vec3f betaTexelSum = 0.0f;
CGU_FLOAT alphaBetaSum = 0.0f;
CGU_FLOAT alphaSqSum = 0.0f;
CGU_FLOAT betaSqSum = 0.0f;
for (CGU_UINT32 i = 0; i < 16; i++)
{
CGU_FLOAT texelPos = cmp_f32tof16(cmp_dotVec3f(texels[i], blockDir));
CGU_UINT32 texelIndex = cmp_Get4BitIndexPos(texelPos, endPoint0Pos, endPoint1Pos);
CGU_FLOAT beta = cmp_saturate(texelIndex / 15.0f);
CGU_FLOAT alpha = 1.0f - beta;
CGU_Vec3f texelF16;
texelF16.x = cmp_f32tof16(texels[i].x);
texelF16.y = cmp_f32tof16(texels[i].y);
texelF16.z = cmp_f32tof16(texels[i].z);
alphaTexelSum += texelF16 * alpha;
betaTexelSum += texelF16 * beta;
alphaBetaSum += alpha * beta;
alphaSqSum += alpha * alpha;
betaSqSum += beta * beta;
}
CGU_FLOAT det = alphaSqSum * betaSqSum - alphaBetaSum * alphaBetaSum;
if (abs(det) > 0.00001f)
{
CGU_FLOAT detRcp = cmp_rcp(det);
blockMin = cmp_clampVec3f((alphaTexelSum * betaSqSum - betaTexelSum * alphaBetaSum) * detRcp, 0.0f, CMP_MAX_16BITFLOAT);
blockMax = cmp_clampVec3f((betaTexelSum * alphaSqSum - alphaTexelSum * alphaBetaSum) * detRcp, 0.0f, CMP_MAX_16BITFLOAT);
blockMin = cmp_f16tof32(blockMin);
blockMax = cmp_f16tof32(blockMax);
}
}
//=============================================================================================
CMP_STATIC CGU_Vec3f cmp_fabsVec3f(CGU_Vec3f value)
{
#ifdef ASPM_HLSL
return abs(value);
#else
CGU_Vec3f res;
res.x = abs(value.x);
res.y = abs(value.y);
res.z = abs(value.z);
return res;
#endif
}
CMP_STATIC CGU_UINT32 cmp_constructColor(CMP_IN CGU_Vec3ui EndPoints)
{
return (((EndPoints.r & 0x000000F8) << 8) | ((EndPoints.g & 0x000000FC) << 3) | ((EndPoints.b & 0x000000F8) >> 3));
}
CMP_STATIC CGU_UINT32 cmp_constructColorBGR(CMP_IN CGU_Vec3f EndPoints)
{
return (((CGU_UINT32(EndPoints.b) & 0x000000F8) << 8) | ((CGU_UINT32(EndPoints.g) & 0x000000FC) << 3) | ((CGU_UINT32(EndPoints.r) & 0x000000F8) >> 3));
}
CMP_STATIC CGU_FLOAT cmp_mod(CMP_IN CGU_FLOAT value, CMP_IN CGU_FLOAT modval)
{
#ifdef ASPM_GLSL
return mod(value, modval);
#endif
return fmod(value, modval);
}
CMP_STATIC CGU_Vec3f cmp_truncVec3f(CMP_IN CGU_Vec3f value)
{
#ifdef ASPM_HLSL
return trunc(value);
#else
CGU_Vec3f res;
res.x = trunc(value.x);
res.y = trunc(value.y);
res.z = trunc(value.z);
return res;
#endif
}
CMP_STATIC CGU_Vec3f cmp_ceilVec3f(CMP_IN CGU_Vec3f value)
{
CGU_Vec3f res;
res.x = ceil(value.x);
res.y = ceil(value.y);
res.z = ceil(value.z);
return res;
}
CMP_STATIC CGU_FLOAT cmp_sqrt(CGU_FLOAT value)
{
return sqrt(value);
}
// Computes inverse square root over an implementation-defined range. The maximum error is implementation-defined.
CMP_STATIC CGV_FLOAT cmp_rsqrt(CGV_FLOAT f)
{
CGV_FLOAT sf = sqrt(f);
if (sf != 0)
return 1 / sqrt(f);
else
return 0.0f;
}
// Common to BC7 API ------------------------------------------------------------------------------------------------------------------------
// valid bit range is 0..8 for mode 1
CMP_STATIC INLINE CGU_UINT32 cmp_shift_right_uint32(CMP_IN CGU_UINT32 v, CMP_IN CGU_INT bits)
{
return v >> bits; // (perf warning expected)
}
CMP_STATIC INLINE CGU_INT cmp_clampi(CMP_IN CGU_INT value, CMP_IN CGU_INT low, CMP_IN CGU_INT high)
{
if (value < low)
return low;
else if (value > high)
return high;
return value;
}
CMP_STATIC INLINE CGU_INT32 cmp_clampi32(CMP_IN CGU_INT32 value, CMP_IN CGU_INT32 low, CMP_IN CGU_INT32 high)
{
if (value < low)
value = low;
else if (value > high)
value = high;
return value;
}
CMP_STATIC CGV_FLOAT cmp_dot4f(CMP_IN CGV_Vec4f value1, CMP_IN CGV_Vec4f value2)
{
#ifdef ASPM_GPU
return dot(value1, value2);
#else
return (value1.x * value2.x) + (value1.y * value2.y) + (value1.z * value2.z) + (value1.w * value2.w);
#endif
}
CMP_STATIC INLINE void cmp_set_vec4f(CMP_INOUT CGU_Vec4f CMP_REFINOUT pV, CMP_IN CGU_FLOAT x, CMP_IN CGU_FLOAT y, CMP_IN CGU_FLOAT z, CMP_IN CGU_FLOAT w)
{
pV[0] = x;
pV[1] = y;
pV[2] = z;
pV[3] = w;
}
CMP_STATIC INLINE void cmp_set_vec4ui(CGU_Vec4ui CMP_REFINOUT pV, CMP_IN CGU_UINT8 x, CMP_IN CGU_UINT8 y, CMP_IN CGU_UINT8 z, CMP_IN CGU_UINT8 w)
{
pV[0] = x;
pV[1] = y;
pV[2] = z;
pV[3] = w;
}
CMP_STATIC inline void cmp_set_vec4ui_clamped(CGU_Vec4ui CMP_REFINOUT pRes, CMP_IN CGU_INT32 r, CMP_IN CGU_INT32 g, CMP_IN CGU_INT32 b, CMP_IN CGU_INT32 a)
{
pRes[0] = (CGU_UINT8)cmp_clampi32(r, 0, 255);
pRes[1] = (CGU_UINT8)cmp_clampi32(g, 0, 255);
pRes[2] = (CGU_UINT8)cmp_clampi32(b, 0, 255);
pRes[3] = (CGU_UINT8)cmp_clampi32(a, 0, 255);
}
CMP_STATIC inline CGU_Vec4f cmp_clampNorm4f(CMP_IN CGU_Vec4f pV)
{
CGU_Vec4f res;
res[0] = cmp_clampf(pV[0], 0.0f, 1.0f);
res[1] = cmp_clampf(pV[1], 0.0f, 1.0f);
res[2] = cmp_clampf(pV[2], 0.0f, 1.0f);
res[3] = cmp_clampf(pV[3], 0.0f, 1.0f);
return res;
}
CMP_STATIC INLINE CGU_Vec4f cmp_vec4ui_to_vec4f(CMP_IN CGU_Vec4ui pC)
{
CGU_Vec4f res;
cmp_set_vec4f(res, (CGU_FLOAT)pC[0], (CGU_FLOAT)pC[1], (CGU_FLOAT)pC[2], (CGU_FLOAT)pC[3]);
return res;
}
CMP_STATIC INLINE void cmp_normalize(CGU_Vec4f CMP_REFINOUT pV)
{
CGU_FLOAT s = cmp_dot4f(pV, pV);
if (s != 0.0f)
{
s = 1.0f / cmp_sqrt(s);
pV *= s;
}
}
CMP_STATIC INLINE CGV_FLOAT cmp_squaref(CMP_IN CGV_FLOAT v)
{
return v * v;
}
CMP_STATIC INLINE CGU_INT cmp_squarei(CMP_IN CGU_INT i)
{
return i * i;
}
CMP_STATIC CGU_UINT8 cmp_clampui8(CMP_IN CGU_UINT8 v, CMP_IN CGU_UINT8 a, CMP_IN CGU_UINT8 b)
{
if (v < a)
return a;
else if (v > b)
return b;
return v;
}
CMP_STATIC CGU_INT32 cmp_abs32(CMP_IN CGU_INT32 v)
{
CGU_UINT32 msk = v >> 31;
return (v ^ msk) - msk;
}
CMP_STATIC void cmp_swap32(CMP_INOUT CGU_UINT32 CMP_REFINOUT a, CMP_INOUT CGU_UINT32 CMP_REFINOUT b)
{
CGU_UINT32 t = a;
a = b;
b = t;
}
// Computes inverse square root over an implementation-defined range. The maximum error is implementation-defined.
CMP_STATIC CGV_FLOAT cmp_Image_rsqrt(CMP_IN CGV_FLOAT f)
{
CGV_FLOAT sf = sqrt(f);
if (sf != 0)
return 1 / sqrt(f);
else
return 0.0f;
}
CMP_STATIC void cmp_pack4bitindex32(CMP_INOUT CGU_UINT32 packed_index[2], CMP_IN CGU_UINT32 src_index[16])
{
// Converts from unpacked index to packed index
packed_index[0] = 0x0000;
packed_index[1] = 0x0000;
CGU_UINT32 shift = 0; // was CGU_UINT8
for (CGU_INT k = 0; k < 8; k++)
{
packed_index[0] |= (CGU_UINT32)(src_index[k] & 0x0F) << shift;
packed_index[1] |= (CGU_UINT32)(src_index[k + 8] & 0x0F) << shift;
shift += 4;
}
}
CMP_STATIC void cmp_pack4bitindex(CMP_INOUT CGU_UINT32 packed_index[2], CMP_IN CGU_UINT8 src_index[16])
{
// Converts from unpacked index to packed index
packed_index[0] = 0x0000;
packed_index[1] = 0x0000;
CGU_UINT32 shift = 0; // was CGU_UINT8
for (CGU_INT k = 0; k < 8; k++)
{
packed_index[0] |= (CGU_UINT32)(src_index[k] & 0x0F) << shift;
packed_index[1] |= (CGU_UINT32)(src_index[k + 8] & 0x0F) << shift;
shift += 4;
}
}
CMP_STATIC INLINE CGU_INT cmp_expandbits(CMP_IN CGU_INT v, CMP_IN CGU_INT bits)
{
CGU_INT vv = v << (8 - bits);
return vv + cmp_shift_right_uint32(vv, bits);
}
// This code need further improvements and investigation
CMP_STATIC INLINE CGU_UINT8 cmp_ep_find_floor2(CMP_IN CGV_FLOAT v, CMP_IN CGU_UINT8 bits, CMP_IN CGU_UINT8 use_par, CMP_IN CGU_UINT8 odd)
{
CGU_UINT8 i1 = 0;
CGU_UINT8 i2 = 1 << (bits - use_par);
odd = use_par ? odd : 0;
while (i2 - i1 > 1)
{
CGU_UINT8 j = (CGU_UINT8)((i1 + i2) * 0.5f);
CGV_FLOAT ep_d = (CGV_FLOAT)cmp_expandbits((j << use_par) + odd, bits);
if (v >= ep_d)
i1 = j;
else
i2 = j;
}
return (i1 << use_par) + odd;
}
CMP_STATIC CGV_FLOAT cmp_absf(CMP_IN CGV_FLOAT a)
{
return a > 0.0F ? a : -a;
}
CMP_STATIC INLINE CGU_UINT32 cmp_pow2Packed(CMP_IN CGU_INT x)
{
return 1 << x;
}
CMP_STATIC INLINE CGU_UINT8 cmp_clampIndex(CMP_IN CGU_UINT8 v, CMP_IN CGU_UINT8 a, CMP_IN CGU_UINT8 b)
{
if (v < a)
return a;
else if (v > b)
return b;
return v;
}
CMP_STATIC INLINE CGU_UINT8 shift_right_uint82(CMP_IN CGU_UINT8 v, CMP_IN CGU_UINT8 bits)
{
return v >> bits; // (perf warning expected)
}
#endif
CMP_STATIC CGU_INT cmp_QuantizeToBitSize(CMP_IN CGU_INT value, CMP_IN CGU_INT prec, CMP_IN CGU_BOOL signedfloat16)
{
if (prec <= 1)
return 0;
CGU_BOOL negvalue = false;
// move data to use extra bits for processing
CGU_INT ivalue = value;
if (signedfloat16)
{
if (value < 0)
{
negvalue = true;
value = -value;
}
prec--;
}
else
{
// clamp -ve
if (value < 0)
value = 0;
}
CGU_INT iQuantized;
CGU_INT bias = (prec > 10 && prec != 16) ? ((1 << (prec - 11)) - 1) : 0;
bias = (prec == 16) ? 15 : bias;
iQuantized = ((ivalue << prec) + bias) / (0x7bff + 1); // 16 bit Float Max 0x7bff
return (negvalue ? -iQuantized : iQuantized);
}
//=======================================================
// CPU GPU Macro API
//=======================================================
#ifdef ASPM_GPU
#define cmp_min(a, b) min(a, b)
#define cmp_max(a, b) max(a, b)
#else
#ifndef cmp_min
#define cmp_min(a, b) ((a) < (b) ? (a) : (b))
#endif
#ifndef cmp_max
#define cmp_max(a, b) ((a) > (b) ? (a) : (b))
#endif
#endif
//=======================================================
// CPU Template API
//=======================================================
#ifndef ASPM_GPU
#ifndef TEMPLATE_API_INTERFACED
#define TEMPLATE_API_INTERFACED
template <typename T>
T clamp(T& v, const T& lo, const T& hi)
{
if (v < lo)
return lo;
else if (v > hi)
return hi;
return v;
}
template <typename T>
Vec4T<T> clamp(Vec4T<T>& v, const T& lo, const T& hi)
{
Vec4T<T> res = v;
if (v.x < lo)
res.x = lo;
else if (v.x > hi)
res.x = hi;
if (v.y < lo)
res.y = lo;
else if (v.y > hi)
res.y = hi;
if (v.z < lo)
res.z = lo;
else if (v.w > hi)
res.w = hi;
if (v.w < lo)
res.w = lo;
else if (v.z > hi)
res.z = hi;
return res;
}
template <typename T,typename T2>
T dot(T& v1, T2& v2)
{
return (v1 * v2);
}
template <typename T>
T dot(Vec4T<T>& v1, Vec4T<T>& v2)
{
return (v1.x * v2.x + v1.y * v2.y + v1.z * v2.z + v1.w * v2.w);
}
template <typename T, typename T2>
T dot(Vec4T<T>& v1, Vec4T<T2>& v2)
{
return (v1.x * v2.x + v1.y * v2.y + v1.z * v2.z + v1.w * v2.w);
}
template <typename T>
T dot(Vec3T<T>& v1, Vec3T<T>& v2)
{
return (v1.x * v2.x + v1.y * v2.y + v1.z * v2.z);
}
template <typename T, typename T2>
T dot(Vec3T<T>& v1, Vec3T<T2>& v2)
{
return (v1.x * v2.x + v1.y * v2.y + v1.z * v2.z);
}
#endif // API_INTERFACED
#endif // ASPM_GPU
#endif //
Reference in New Issue View Git Blame Copy Permalink