1341 lines
36 KiB
C++
1341 lines
36 KiB
C++
#pragma once
|
|
|
|
#include <stdint.h>
|
|
|
|
namespace cacheflow {
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
struct Float8_ {
|
|
float2 x;
|
|
float2 y;
|
|
float2 z;
|
|
float2 w;
|
|
};
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
struct Float4_ {
|
|
float2 x;
|
|
float2 y;
|
|
};
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
#ifdef ENABLE_BF16
|
|
struct bf16_4_t {
|
|
__nv_bfloat162 x;
|
|
__nv_bfloat162 y;
|
|
};
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
struct bf16_8_t {
|
|
__nv_bfloat162 x;
|
|
__nv_bfloat162 y;
|
|
__nv_bfloat162 z;
|
|
__nv_bfloat162 w;
|
|
};
|
|
#endif
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float add(float a, float b)
|
|
{
|
|
return a + b;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 add(float2 a, float2 b)
|
|
{
|
|
float2 c;
|
|
c.x = add(a.x, b.x);
|
|
c.y = add(a.y, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float4 add(float4 a, float4 b)
|
|
{
|
|
float4 c;
|
|
c.x = add(a.x, b.x);
|
|
c.y = add(a.y, b.y);
|
|
c.z = add(a.z, b.z);
|
|
c.w = add(a.w, b.w);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
#ifdef ENABLE_BF16
|
|
inline __device__ __nv_bfloat16 add(__nv_bfloat16 a, __nv_bfloat16 b)
|
|
{
|
|
return a + b;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ __nv_bfloat162 add(__nv_bfloat162 a, __nv_bfloat162 b)
|
|
{
|
|
return bf16hadd2(a, b);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ bf16_4_t add(bf16_4_t a, bf16_4_t b)
|
|
{
|
|
bf16_4_t c;
|
|
c.x = add(a.x, b.x);
|
|
c.y = add(a.y, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ bf16_8_t add(bf16_8_t a, bf16_8_t b)
|
|
{
|
|
bf16_8_t c;
|
|
c.x = add(a.x, b.x);
|
|
c.y = add(a.y, b.y);
|
|
c.z = add(a.z, b.z);
|
|
c.w = add(a.w, b.w);
|
|
return c;
|
|
}
|
|
#endif // ENABLE_BF16
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint16_t add(uint16_t a, uint16_t b)
|
|
{
|
|
uint16_t c;
|
|
asm volatile("add.f16 %0, %1, %2;\n" : "=h"(c) : "h"(a), "h"(b));
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint32_t add(uint32_t a, uint32_t b)
|
|
{
|
|
uint32_t c;
|
|
asm volatile("add.f16x2 %0, %1, %2;\n" : "=r"(c) : "r"(a), "r"(b));
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint2 add(uint2 a, uint2 b)
|
|
{
|
|
uint2 c;
|
|
c.x = add(a.x, b.x);
|
|
c.y = add(a.y, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint4 add(uint4 a, uint4 b)
|
|
{
|
|
uint4 c;
|
|
c.x = add(a.x, b.x);
|
|
c.y = add(a.y, b.y);
|
|
c.z = add(a.z, b.z);
|
|
c.w = add(a.w, b.w);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint16_t float_to_half(float f)
|
|
{
|
|
union {
|
|
uint32_t u32;
|
|
uint16_t u16[2];
|
|
} tmp;
|
|
#if 0 && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 // Is it better?
|
|
float zero = 0.f;
|
|
asm volatile("cvt.rn.f16x2.f32 %0, %1, %2;\n" : "=r"(tmp.u32) : "f"(zero), "f"(f));
|
|
#else
|
|
asm volatile("cvt.rn.f16.f32 %0, %1;\n" : "=h"(tmp.u16[0]) : "f"(f));
|
|
#endif
|
|
return tmp.u16[0];
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint32_t float2_to_half2(float2 f)
|
|
{
|
|
union {
|
|
uint32_t u32;
|
|
uint16_t u16[2];
|
|
} tmp;
|
|
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
|
|
asm volatile("cvt.rn.f16x2.f32 %0, %1, %2;\n" : "=r"(tmp.u32) : "f"(f.y), "f"(f.x));
|
|
#else
|
|
asm volatile("cvt.rn.f16.f32 %0, %1;\n" : "=h"(tmp.u16[0]) : "f"(f.x));
|
|
asm volatile("cvt.rn.f16.f32 %0, %1;\n" : "=h"(tmp.u16[1]) : "f"(f.y));
|
|
#endif
|
|
return tmp.u32;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float half_to_float(uint16_t h)
|
|
{
|
|
float f;
|
|
asm volatile("cvt.f32.f16 %0, %1;\n" : "=f"(f) : "h"(h));
|
|
return f;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 half2_to_float2(uint32_t v)
|
|
{
|
|
uint16_t lo, hi;
|
|
asm volatile("mov.b32 {%0, %1}, %2;\n" : "=h"(lo), "=h"(hi) : "r"(v));
|
|
return make_float2(half_to_float(lo), half_to_float(hi));
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 add(uint32_t a, float2 fb)
|
|
{
|
|
float2 fa = half2_to_float2(a);
|
|
return add(fa, fb);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ add(uint2 a, Float4_ fb)
|
|
{
|
|
Float4_ fc;
|
|
fc.x = add(a.x, fb.x);
|
|
fc.y = add(a.y, fb.y);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ add(uint4 a, Float8_ fb)
|
|
{
|
|
Float8_ fc;
|
|
fc.x = add(a.x, fb.x);
|
|
fc.y = add(a.y, fb.y);
|
|
fc.z = add(a.z, fb.z);
|
|
fc.w = add(a.w, fb.w);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint32_t h0_h0(uint16_t a)
|
|
{
|
|
uint32_t b;
|
|
asm volatile("mov.b32 %0, {%1, %1};" : "=r"(b) : "h"(a));
|
|
return b;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float fma(float a, float b, float c)
|
|
{
|
|
return a * b + c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 fma(float2 a, float2 b, float2 c)
|
|
{
|
|
float2 d;
|
|
d.x = fma(a.x, b.x, c.x);
|
|
d.y = fma(a.y, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 fma(float a, float2 b, float2 c)
|
|
{
|
|
float2 d;
|
|
d.x = fma(a, b.x, c.x);
|
|
d.y = fma(a, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float4 fma(float4 a, float4 b, float4 c)
|
|
{
|
|
float4 d;
|
|
d.x = fma(a.x, b.x, c.x);
|
|
d.y = fma(a.y, b.y, c.y);
|
|
d.z = fma(a.z, b.z, c.z);
|
|
d.w = fma(a.w, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float4 fma(float a, float4 b, float4 c)
|
|
{
|
|
float4 d;
|
|
d.x = fma(a, b.x, c.x);
|
|
d.y = fma(a, b.y, c.y);
|
|
d.z = fma(a, b.z, c.z);
|
|
d.w = fma(a, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ fma(float a, Float4_ b, Float4_ c)
|
|
{
|
|
Float4_ d;
|
|
d.x = fma(a, b.x, c.x);
|
|
d.y = fma(a, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ fma(float a, Float8_ b, Float8_ c)
|
|
{
|
|
Float8_ d;
|
|
d.x = fma(a, b.x, c.x);
|
|
d.y = fma(a, b.y, c.y);
|
|
d.z = fma(a, b.z, c.z);
|
|
d.w = fma(a, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
#ifdef ENABLE_BF16
|
|
inline __device__ float2 add(__nv_bfloat162 a, float2 fb)
|
|
{
|
|
float2 fa = bf1622float2(a);
|
|
return add(fa, fb);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ add(bf16_4_t a, Float4_ fb)
|
|
{
|
|
Float4_ fc;
|
|
fc.x = add(a.x, fb.x);
|
|
fc.y = add(a.y, fb.y);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ add(bf16_8_t a, Float8_ fb)
|
|
{
|
|
Float8_ fc;
|
|
fc.x = add(a.x, fb.x);
|
|
fc.y = add(a.y, fb.y);
|
|
fc.z = add(a.z, fb.z);
|
|
fc.w = add(a.w, fb.w);
|
|
return fc;
|
|
}
|
|
#endif // ENABLE_BF16
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint32_t fma(uint32_t a, uint32_t b, uint32_t c)
|
|
{
|
|
uint32_t d;
|
|
asm volatile("fma.rn.f16x2 %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "r"(b), "r"(c));
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint32_t fma(uint16_t a, uint32_t b, uint32_t c)
|
|
{
|
|
return fma(h0_h0(a), b, c);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint2 fma(uint2 a, uint2 b, uint2 c)
|
|
{
|
|
uint2 d;
|
|
d.x = fma(a.x, b.x, c.x);
|
|
d.y = fma(a.y, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint2 fma(uint16_t a, uint2 b, uint2 c)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
uint2 d;
|
|
d.x = fma(s, b.x, c.x);
|
|
d.y = fma(s, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint4 fma(uint4 a, uint4 b, uint4 c)
|
|
{
|
|
uint4 d;
|
|
d.x = fma(a.x, b.x, c.x);
|
|
d.y = fma(a.y, b.y, c.y);
|
|
d.z = fma(a.z, b.z, c.z);
|
|
d.w = fma(a.w, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ uint4 fma(uint16_t a, uint4 b, uint4 c)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
uint4 d;
|
|
d.x = fma(s, b.x, c.x);
|
|
d.y = fma(s, b.y, c.y);
|
|
d.z = fma(s, b.z, c.z);
|
|
d.w = fma(s, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float fma(uint16_t a, uint16_t b, float fc)
|
|
{
|
|
float fa = half_to_float(a);
|
|
float fb = half_to_float(b);
|
|
return fa * fb + fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 fma(uint32_t a, uint32_t b, float2 fc)
|
|
{
|
|
float2 fa = half2_to_float2(a);
|
|
float2 fb = half2_to_float2(b);
|
|
return fma(fa, fb, fc);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 fma(uint16_t a, uint32_t b, float2 fc)
|
|
{
|
|
return fma(h0_h0(a), b, fc);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ fma(uint2 a, uint2 b, Float4_ fc)
|
|
{
|
|
Float4_ fd;
|
|
fd.x = fma(a.x, b.x, fc.x);
|
|
fd.y = fma(a.y, b.y, fc.y);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ fma(uint16_t a, uint2 b, Float4_ fc)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
Float4_ fd;
|
|
fd.x = fma(s, b.x, fc.x);
|
|
fd.y = fma(s, b.y, fc.y);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ fma(uint4 a, uint4 b, Float8_ fc)
|
|
{
|
|
Float8_ fd;
|
|
fd.x = fma(a.x, b.x, fc.x);
|
|
fd.y = fma(a.y, b.y, fc.y);
|
|
fd.z = fma(a.z, b.z, fc.z);
|
|
fd.w = fma(a.w, b.w, fc.w);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ fma(uint16_t a, uint4 b, Float8_ fc)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
Float8_ fd;
|
|
fd.x = fma(s, b.x, fc.x);
|
|
fd.y = fma(s, b.y, fc.y);
|
|
fd.z = fma(s, b.z, fc.z);
|
|
fd.w = fma(s, b.w, fc.w);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
#ifdef ENABLE_BF16
|
|
inline __device__ __nv_bfloat162 fma(__nv_bfloat162 a, __nv_bfloat162 b, __nv_bfloat162 c)
|
|
{
|
|
return bf16hfma2(a, b, c);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ __nv_bfloat162 fma(__nv_bfloat16 a, __nv_bfloat162 b, __nv_bfloat162 c)
|
|
{
|
|
return bf16hfma2(bf162bf162(a), b, c);
|
|
}
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ bf16_4_t fma(bf16_4_t a, bf16_4_t b, bf16_4_t c)
|
|
{
|
|
bf16_4_t d;
|
|
d.x = fma(a.x, b.x, c.x);
|
|
d.y = fma(a.y, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ bf16_4_t fma(__nv_bfloat16 a, bf16_4_t b, bf16_4_t c)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
bf16_4_t d;
|
|
d.x = fma(s, b.x, c.x);
|
|
d.y = fma(s, b.y, c.y);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ bf16_8_t fma(bf16_8_t a, bf16_8_t b, bf16_8_t c)
|
|
{
|
|
bf16_8_t d;
|
|
d.x = fma(a.x, b.x, c.x);
|
|
d.y = fma(a.y, b.y, c.y);
|
|
d.z = fma(a.z, b.z, c.z);
|
|
d.w = fma(a.w, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ bf16_8_t fma(__nv_bfloat16 a, bf16_8_t b, bf16_8_t c)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
bf16_8_t d;
|
|
d.x = fma(s, b.x, c.x);
|
|
d.y = fma(s, b.y, c.y);
|
|
d.z = fma(s, b.z, c.z);
|
|
d.w = fma(s, b.w, c.w);
|
|
return d;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float fma(__nv_bfloat16 a, __nv_bfloat16 b, float fc)
|
|
{
|
|
return __bfloat162float(a) * __bfloat162float(b) + fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 fma(__nv_bfloat162 a, __nv_bfloat162 b, float2 fc)
|
|
{
|
|
float2 fa = bf1622float2(a);
|
|
float2 fb = bf1622float2(b);
|
|
return fma(fa, fb, fc);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 fma(__nv_bfloat16 a, __nv_bfloat162 b, float2 fc)
|
|
{
|
|
return fma(bf162bf162(a), b, fc);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ fma(bf16_4_t a, bf16_4_t b, Float4_ fc)
|
|
{
|
|
Float4_ fd;
|
|
fd.x = fma(a.x, b.x, fc.x);
|
|
fd.y = fma(a.y, b.y, fc.y);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ fma(__nv_bfloat16 a, bf16_4_t b, Float4_ fc)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
Float4_ fd;
|
|
fd.x = fma(s, b.x, fc.x);
|
|
fd.y = fma(s, b.y, fc.y);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ fma(bf16_8_t a, bf16_8_t b, Float8_ fc)
|
|
{
|
|
Float8_ fd;
|
|
fd.x = fma(a.x, b.x, fc.x);
|
|
fd.y = fma(a.y, b.y, fc.y);
|
|
fd.z = fma(a.z, b.z, fc.z);
|
|
fd.w = fma(a.w, b.w, fc.w);
|
|
return fd;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ fma(__nv_bfloat16 a, bf16_8_t b, Float8_ fc)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
Float8_ fd;
|
|
fd.x = fma(s, b.x, fc.x);
|
|
fd.y = fma(s, b.y, fc.y);
|
|
fd.z = fma(s, b.z, fc.z);
|
|
fd.w = fma(s, b.w, fc.w);
|
|
return fd;
|
|
}
|
|
#endif // ENABLE_BF16
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<typename Acc, typename A, typename B>
|
|
inline __device__ Acc mul(A a, B b);
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float mul<float, float>(float a, float b)
|
|
{
|
|
return a * b;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float2 mul(float2 a, float2 b)
|
|
{
|
|
float2 c;
|
|
c.x = a.x * b.x;
|
|
c.y = a.y * b.y;
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float2 mul(float a, float2 b)
|
|
{
|
|
float2 c;
|
|
c.x = a * b.x;
|
|
c.y = a * b.y;
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float4 mul(float4 a, float4 b)
|
|
{
|
|
float4 c;
|
|
c.x = a.x * b.x;
|
|
c.y = a.y * b.y;
|
|
c.z = a.z * b.z;
|
|
c.w = a.w * b.w;
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float4 mul(float a, float4 b)
|
|
{
|
|
float4 c;
|
|
c.x = a * b.x;
|
|
c.y = a * b.y;
|
|
c.z = a * b.z;
|
|
c.w = a * b.w;
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint16_t mul(uint16_t a, uint16_t b)
|
|
{
|
|
uint16_t c;
|
|
asm volatile("mul.f16 %0, %1, %2;\n" : "=h"(c) : "h"(a), "h"(b));
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint32_t mul(uint32_t a, uint32_t b)
|
|
{
|
|
uint32_t c;
|
|
asm volatile("mul.f16x2 %0, %1, %2;\n" : "=r"(c) : "r"(a), "r"(b));
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint32_t mul(uint16_t a, uint32_t b)
|
|
{
|
|
return mul<uint32_t, uint32_t, uint32_t>(h0_h0(a), b);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint2 mul(uint2 a, uint2 b)
|
|
{
|
|
uint2 c;
|
|
c.x = mul<uint32_t, uint32_t, uint32_t>(a.x, b.x);
|
|
c.y = mul<uint32_t, uint32_t, uint32_t>(a.y, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint2 mul(uint16_t a, uint2 b)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
uint2 c;
|
|
c.x = mul<uint32_t, uint32_t, uint32_t>(s, b.x);
|
|
c.y = mul<uint32_t, uint32_t, uint32_t>(s, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint4 mul(uint4 a, uint4 b)
|
|
{
|
|
uint4 c;
|
|
c.x = mul<uint32_t, uint32_t, uint32_t>(a.x, b.x);
|
|
c.y = mul<uint32_t, uint32_t, uint32_t>(a.y, b.y);
|
|
c.z = mul<uint32_t, uint32_t, uint32_t>(a.z, b.z);
|
|
c.w = mul<uint32_t, uint32_t, uint32_t>(a.w, b.w);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ uint4 mul(uint16_t a, uint4 b)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
uint4 c;
|
|
c.x = mul<uint32_t, uint32_t, uint32_t>(s, b.x);
|
|
c.y = mul<uint32_t, uint32_t, uint32_t>(s, b.y);
|
|
c.z = mul<uint32_t, uint32_t, uint32_t>(s, b.z);
|
|
c.w = mul<uint32_t, uint32_t, uint32_t>(s, b.w);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float mul(uint16_t a, uint16_t b)
|
|
{
|
|
float fa = half_to_float(a);
|
|
float fb = half_to_float(b);
|
|
return fa * fb;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float2 mul(uint32_t a, uint32_t b)
|
|
{
|
|
float2 fa = half2_to_float2(a);
|
|
float2 fb = half2_to_float2(b);
|
|
return mul<float2, float2, float2>(fa, fb);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float2 mul(uint16_t a, uint32_t b)
|
|
{
|
|
return mul<float2, uint32_t, uint32_t>(h0_h0(a), b);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float4_ mul(uint2 a, uint2 b)
|
|
{
|
|
Float4_ fc;
|
|
fc.x = mul<float2, uint32_t, uint32_t>(a.x, b.x);
|
|
fc.y = mul<float2, uint32_t, uint32_t>(a.y, b.y);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float4_ mul(uint16_t a, uint2 b)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
Float4_ fc;
|
|
fc.x = mul<float2, uint32_t, uint32_t>(s, b.x);
|
|
fc.y = mul<float2, uint32_t, uint32_t>(s, b.y);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float8_ mul(uint4 a, uint4 b)
|
|
{
|
|
Float8_ fc;
|
|
fc.x = mul<float2, uint32_t, uint32_t>(a.x, b.x);
|
|
fc.y = mul<float2, uint32_t, uint32_t>(a.y, b.y);
|
|
fc.z = mul<float2, uint32_t, uint32_t>(a.z, b.z);
|
|
fc.w = mul<float2, uint32_t, uint32_t>(a.w, b.w);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float8_ mul(uint16_t a, uint4 b)
|
|
{
|
|
uint32_t s = h0_h0(a);
|
|
Float8_ fc;
|
|
fc.x = mul<float2, uint32_t, uint32_t>(s, b.x);
|
|
fc.y = mul<float2, uint32_t, uint32_t>(s, b.y);
|
|
fc.z = mul<float2, uint32_t, uint32_t>(s, b.z);
|
|
fc.w = mul<float2, uint32_t, uint32_t>(s, b.w);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
#ifdef ENABLE_BF16
|
|
template<>
|
|
inline __device__ __nv_bfloat16 mul(__nv_bfloat16 a, __nv_bfloat16 b)
|
|
{
|
|
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
|
|
return __hmul(a, b);
|
|
#else
|
|
return bf16hmul(a, b);
|
|
#endif
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ __nv_bfloat162 mul(__nv_bfloat162 a, __nv_bfloat162 b)
|
|
{
|
|
return bf16hmul2(a, b);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ __nv_bfloat162 mul(__nv_bfloat16 a, __nv_bfloat162 b)
|
|
{
|
|
return mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(bf162bf162(a), b);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ bf16_4_t mul(bf16_4_t a, bf16_4_t b)
|
|
{
|
|
bf16_4_t c;
|
|
c.x = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(a.x, b.x);
|
|
c.y = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(a.y, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ bf16_4_t mul(__nv_bfloat16 a, bf16_4_t b)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
bf16_4_t c;
|
|
c.x = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(s, b.x);
|
|
c.y = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(s, b.y);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ bf16_8_t mul(bf16_8_t a, bf16_8_t b)
|
|
{
|
|
bf16_8_t c;
|
|
c.x = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(a.x, b.x);
|
|
c.y = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(a.y, b.y);
|
|
c.z = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(a.z, b.z);
|
|
c.w = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(a.w, b.w);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ bf16_8_t mul(__nv_bfloat16 a, bf16_8_t b)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
bf16_8_t c;
|
|
c.x = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(s, b.x);
|
|
c.y = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(s, b.y);
|
|
c.z = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(s, b.z);
|
|
c.w = mul<__nv_bfloat162, __nv_bfloat162, __nv_bfloat162>(s, b.w);
|
|
return c;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float mul(__nv_bfloat16 a, __nv_bfloat16 b)
|
|
{
|
|
float fa = (float)a;
|
|
float fb = (float)b;
|
|
return fa * fb;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float2 mul(__nv_bfloat162 a, __nv_bfloat162 b)
|
|
{
|
|
float2 fa = bf1622float2(a);
|
|
float2 fb = bf1622float2(b);
|
|
return mul<float2, float2, float2>(fa, fb);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ float2 mul(__nv_bfloat16 a, __nv_bfloat162 b)
|
|
{
|
|
return mul<float2, __nv_bfloat162, __nv_bfloat162>(bf162bf162(a), b);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float4_ mul(bf16_4_t a, bf16_4_t b)
|
|
{
|
|
Float4_ fc;
|
|
fc.x = mul<float2, __nv_bfloat162, __nv_bfloat162>(a.x, b.x);
|
|
fc.y = mul<float2, __nv_bfloat162, __nv_bfloat162>(a.y, b.y);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float4_ mul(__nv_bfloat16 a, bf16_4_t b)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
Float4_ fc;
|
|
fc.x = mul<float2, __nv_bfloat162, __nv_bfloat162>(s, b.x);
|
|
fc.y = mul<float2, __nv_bfloat162, __nv_bfloat162>(s, b.y);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float8_ mul(bf16_8_t a, bf16_8_t b)
|
|
{
|
|
Float8_ fc;
|
|
fc.x = mul<float2, __nv_bfloat162, __nv_bfloat162>(a.x, b.x);
|
|
fc.y = mul<float2, __nv_bfloat162, __nv_bfloat162>(a.y, b.y);
|
|
fc.z = mul<float2, __nv_bfloat162, __nv_bfloat162>(a.z, b.z);
|
|
fc.w = mul<float2, __nv_bfloat162, __nv_bfloat162>(a.w, b.w);
|
|
return fc;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<>
|
|
inline __device__ Float8_ mul(__nv_bfloat16 a, bf16_8_t b)
|
|
{
|
|
__nv_bfloat162 s = bf162bf162(a);
|
|
Float8_ fc;
|
|
fc.x = mul<float2, __nv_bfloat162, __nv_bfloat162>(s, b.x);
|
|
fc.y = mul<float2, __nv_bfloat162, __nv_bfloat162>(s, b.y);
|
|
fc.z = mul<float2, __nv_bfloat162, __nv_bfloat162>(s, b.z);
|
|
fc.w = mul<float2, __nv_bfloat162, __nv_bfloat162>(s, b.w);
|
|
return fc;
|
|
}
|
|
#endif // ENABLE_BF16
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(float v)
|
|
{
|
|
return v;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(float2 v)
|
|
{
|
|
return v.x + v.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(float4 v)
|
|
{
|
|
return v.x + v.y + v.z + v.w;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
#ifdef ENABLE_BF16
|
|
inline __device__ float sum(__nv_bfloat162 v)
|
|
{
|
|
float2 vf = bf1622float2(v);
|
|
return vf.x + vf.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(bf16_4_t v)
|
|
{
|
|
return sum(v.x) + sum(v.y);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(bf16_8_t v)
|
|
{
|
|
return sum(v.x) + sum(v.y) + sum(v.z) + sum(v.w);
|
|
}
|
|
#endif // ENABLE_BF16
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(uint16_t v)
|
|
{
|
|
return half_to_float(v);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(uint32_t v)
|
|
{
|
|
float2 tmp = half2_to_float2(v);
|
|
return tmp.x + tmp.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(uint2 v)
|
|
{
|
|
uint32_t c = add(v.x, v.y);
|
|
return sum(c);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(uint4 v)
|
|
{
|
|
#if 1
|
|
uint32_t c = add(v.x, v.y);
|
|
c = add(c, v.z);
|
|
c = add(c, v.w);
|
|
#else
|
|
uint32_t c = add(v.x, v.y);
|
|
uint32_t d = add(v.z, v.w);
|
|
c = add(c, d);
|
|
#endif
|
|
return sum(c);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(Float4_ v)
|
|
{
|
|
return v.x.x + v.x.y + v.y.x + v.y.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float sum(Float8_ v)
|
|
{
|
|
return v.x.x + v.x.y + v.y.x + v.y.y + v.z.x + v.z.y + v.w.x + v.w.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float dot(float a, float b)
|
|
{
|
|
return a * b;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float dot(float2 a, float2 b)
|
|
{
|
|
float2 c = mul<float2, float2, float2>(a, b);
|
|
return c.x + c.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float dot(Float4_ a, Float4_ b)
|
|
{
|
|
float2 acc = mul<float2, float2, float2>(a.x, b.x);
|
|
acc = fma(a.y, b.y, acc);
|
|
return acc.x + acc.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float dot(Float8_ a, Float8_ b)
|
|
{
|
|
float2 acc = mul<float2, float2, float2>(a.x, b.x);
|
|
acc = fma(a.y, b.y, acc);
|
|
acc = fma(a.z, b.z, acc);
|
|
acc = fma(a.w, b.w, acc);
|
|
return acc.x + acc.y;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<typename T>
|
|
inline __device__ float dot(T a, T b)
|
|
{
|
|
return sum(mul<T, T, T>(a, b));
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<typename A, typename T>
|
|
inline __device__ float dot(T a, T b)
|
|
{
|
|
return sum(mul<A, T, T>(a, b));
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void zero(uint16_t& dst)
|
|
{
|
|
dst = uint16_t(0);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
template<typename T>
|
|
inline __device__ void zero(T& dst)
|
|
{
|
|
constexpr int WORDS = sizeof(T) / 4;
|
|
union {
|
|
T raw;
|
|
uint32_t words[WORDS];
|
|
} tmp;
|
|
#pragma unroll
|
|
for (int ii = 0; ii < WORDS; ++ii) {
|
|
tmp.words[ii] = 0u;
|
|
}
|
|
dst = tmp.raw;
|
|
}
|
|
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(float& dst, float src)
|
|
{
|
|
dst = src;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(uint16_t& dst, float src)
|
|
{
|
|
dst = float_to_half(src);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(uint32_t& dst, float2 src)
|
|
{
|
|
dst = float2_to_half2(src);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
#ifdef ENABLE_BF16
|
|
inline __device__ void convert_from_float(__nv_bfloat16& dst, float src)
|
|
{
|
|
dst = __float2bfloat16(src);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(__nv_bfloat162& dst, float2 src)
|
|
{
|
|
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
|
|
dst = __float22bfloat162_rn(src);
|
|
#else
|
|
dst = __floats2bfloat162_rn(src.x, src.y);
|
|
#endif
|
|
}
|
|
#endif // ENABLE_BF16
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(uint2& dst, Float4_ src)
|
|
{
|
|
dst.x = float2_to_half2(src.x);
|
|
dst.y = float2_to_half2(src.y);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(uint4& dst, Float8_ src)
|
|
{
|
|
dst.x = float2_to_half2(src.x);
|
|
dst.y = float2_to_half2(src.y);
|
|
dst.z = float2_to_half2(src.z);
|
|
dst.w = float2_to_half2(src.w);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
#ifdef ENABLE_BF16
|
|
inline __device__ void convert_from_float(bf16_4_t& dst, Float4_ src)
|
|
{
|
|
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
|
|
dst.x = __float22bfloat162_rn(src.x);
|
|
dst.y = __float22bfloat162_rn(src.y);
|
|
#else
|
|
dst.x = __floats2bfloat162_rn(src.x.x, src.x.y);
|
|
dst.y = __floats2bfloat162_rn(src.y.x, src.y.y);
|
|
#endif
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(bf16_8_t& dst, Float8_ src)
|
|
{
|
|
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
|
|
dst.x = __float22bfloat162_rn(src.x);
|
|
dst.y = __float22bfloat162_rn(src.y);
|
|
dst.z = __float22bfloat162_rn(src.z);
|
|
dst.w = __float22bfloat162_rn(src.w);
|
|
#else
|
|
dst.x = __floats2bfloat162_rn(src.x.x, src.x.y);
|
|
dst.y = __floats2bfloat162_rn(src.y.x, src.y.y);
|
|
dst.z = __floats2bfloat162_rn(src.z.x, src.z.y);
|
|
dst.w = __floats2bfloat162_rn(src.w.x, src.w.y);
|
|
#endif
|
|
}
|
|
#endif // ENABLE_BF16
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(float2& dst, float2 src)
|
|
{
|
|
dst = src;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ void convert_from_float(float4& dst, float4 src)
|
|
{
|
|
dst = src;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float convert_to_float(float4 u)
|
|
{
|
|
return u.x;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float convert_to_float(uint4 u)
|
|
{
|
|
float2 tmp = half2_to_float2(u.x);
|
|
return tmp.x;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
// inline __device__ float cast_to_float(float u)
|
|
// {
|
|
// return u;
|
|
// }
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
// inline __device__ float2 cast_to_float(float2 u)
|
|
// {
|
|
// return u;
|
|
// }
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
// inline __device__ float4 cast_to_float(float4 u)
|
|
// {
|
|
// return u;
|
|
// }
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
// inline __device__ Float4_ cast_to_float(Float4_ u)
|
|
// {
|
|
// return u;
|
|
// }
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
// inline __device__ Float8_ cast_to_float(Float8_ u)
|
|
// {
|
|
// return u;
|
|
// }
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float cast_to_float(uint16_t u)
|
|
{
|
|
return half_to_float(u);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ float2 cast_to_float(uint32_t u)
|
|
{
|
|
return half2_to_float2(u);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float4_ cast_to_float(uint2 u)
|
|
{
|
|
Float4_ tmp;
|
|
tmp.x = half2_to_float2(u.x);
|
|
tmp.y = half2_to_float2(u.y);
|
|
return tmp;
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////////////////////////
|
|
|
|
inline __device__ Float8_ cast_to_float(uint4 u)
|
|
{
|
|
Float8_ tmp;
|
|
tmp.x = half2_to_float2(u.x);
|
|
tmp.y = half2_to_float2(u.y);
|
|
tmp.z = half2_to_float2(u.z);
|
|
tmp.w = half2_to_float2(u.w);
|
|
return tmp;
|
|
}
|
|
|
|
}
|