|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
#include <c10/util/Half.h> |
|
#include "filtered_lrelu.h" |
|
#include <cstdint> |
|
|
|
|
|
|
|
|
|
enum |
|
{ |
|
MODE_SUSD = 0, |
|
MODE_FUSD = 1, |
|
MODE_SUFD = 2, |
|
MODE_FUFD = 3, |
|
}; |
|
|
|
template <class T> struct InternalType; |
|
template <> struct InternalType<double> |
|
{ |
|
typedef double scalar_t; typedef double2 vec2_t; typedef double4 vec4_t; |
|
__device__ __forceinline__ static vec2_t zero_vec2(void) { return make_double2(0, 0); } |
|
__device__ __forceinline__ static vec4_t zero_vec4(void) { return make_double4(0, 0, 0, 0); } |
|
__device__ __forceinline__ static double clamp(double x, double c) { return fmin(fmax(x, -c), c); } |
|
}; |
|
template <> struct InternalType<float> |
|
{ |
|
typedef float scalar_t; typedef float2 vec2_t; typedef float4 vec4_t; |
|
__device__ __forceinline__ static vec2_t zero_vec2(void) { return make_float2(0, 0); } |
|
__device__ __forceinline__ static vec4_t zero_vec4(void) { return make_float4(0, 0, 0, 0); } |
|
__device__ __forceinline__ static float clamp(float x, float c) { return fminf(fmaxf(x, -c), c); } |
|
}; |
|
template <> struct InternalType<c10::Half> |
|
{ |
|
typedef float scalar_t; typedef float2 vec2_t; typedef float4 vec4_t; |
|
__device__ __forceinline__ static vec2_t zero_vec2(void) { return make_float2(0, 0); } |
|
__device__ __forceinline__ static vec4_t zero_vec4(void) { return make_float4(0, 0, 0, 0); } |
|
__device__ __forceinline__ static float clamp(float x, float c) { return fminf(fmaxf(x, -c), c); } |
|
}; |
|
|
|
#define MIN(A, B) ((A) < (B) ? (A) : (B)) |
|
#define MAX(A, B) ((A) > (B) ? (A) : (B)) |
|
#define CEIL_DIV(A, B) (((B)==1) ? (A) : \ |
|
((B)==2) ? ((int)((A)+1) >> 1) : \ |
|
((B)==4) ? ((int)((A)+3) >> 2) : \ |
|
(((A) + ((A) > 0 ? (B) - 1 : 0)) / (B))) |
|
|
|
|
|
template <int N> __device__ __forceinline__ void fast_div_mod(int& x, int& y, unsigned int i) |
|
{ |
|
if ((N & (N-1)) && N <= 256) |
|
y = (i * ((1<<24)/N + 1)) >> 24; |
|
else |
|
y = i/N; |
|
|
|
x = i - y*N; |
|
} |
|
|
|
|
|
template <class T> __device__ __forceinline__ T get_stride(const int64_t& x) |
|
{ |
|
return *reinterpret_cast<const T*>(&x); |
|
} |
|
|
|
|
|
|
|
|
|
#define MAX_FILTER_SIZE 32 |
|
|
|
|
|
__device__ float g_fbuf[2 * MAX_FILTER_SIZE * MAX_FILTER_SIZE]; |
|
__device__ __constant__ float c_fbuf[2 * MAX_FILTER_SIZE * MAX_FILTER_SIZE]; |
|
|
|
|
|
#define c_fu (c_fbuf) |
|
#define c_fd (c_fbuf + MAX_FILTER_SIZE * MAX_FILTER_SIZE) |
|
#define g_fu (g_fbuf) |
|
#define g_fd (g_fbuf + MAX_FILTER_SIZE * MAX_FILTER_SIZE) |
|
|
|
|
|
static __global__ void setup_filters_kernel(filtered_lrelu_kernel_params p) |
|
{ |
|
for (int idx = threadIdx.x; idx < MAX_FILTER_SIZE * MAX_FILTER_SIZE; idx += blockDim.x) |
|
{ |
|
int x, y; |
|
fast_div_mod<MAX_FILTER_SIZE>(x, y, idx); |
|
|
|
int fu_x = p.flip ? x : (p.fuShape.x - 1 - x); |
|
int fu_y = p.flip ? y : (p.fuShape.y - 1 - y); |
|
if (p.fuShape.y > 0) |
|
g_fu[idx] = (x >= p.fuShape.x || y >= p.fuShape.y) ? 0.0f : p.fu[fu_x * p.fuStride.x + fu_y * p.fuStride.y]; |
|
else |
|
g_fu[idx] = (x >= p.fuShape.x || y > 0) ? 0.0f : p.fu[fu_x * p.fuStride.x]; |
|
|
|
int fd_x = p.flip ? x : (p.fdShape.x - 1 - x); |
|
int fd_y = p.flip ? y : (p.fdShape.y - 1 - y); |
|
if (p.fdShape.y > 0) |
|
g_fd[idx] = (x >= p.fdShape.x || y >= p.fdShape.y) ? 0.0f : p.fd[fd_x * p.fdStride.x + fd_y * p.fdStride.y]; |
|
else |
|
g_fd[idx] = (x >= p.fdShape.x || y > 0) ? 0.0f : p.fd[fd_x * p.fdStride.x]; |
|
} |
|
} |
|
|
|
|
|
template <bool, bool> static cudaError_t copy_filters(cudaStream_t stream) |
|
{ |
|
void* src = 0; |
|
cudaError_t err = cudaGetSymbolAddress(&src, g_fbuf); |
|
if (err) return err; |
|
return cudaMemcpyToSymbolAsync(c_fbuf, src, 2 * MAX_FILTER_SIZE * MAX_FILTER_SIZE * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream); |
|
} |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
extern __shared__ char s_buf_raw[]; |
|
|
|
template <class T, class index_t, int sharedKB, bool signWrite, bool signRead, int filterMode, int up, int fuSize, int down, int fdSize, int tileOutW, int tileOutH, int threadsPerBlock, bool enableXrep, bool enableWriteSkip> |
|
static __global__ void filtered_lrelu_kernel(filtered_lrelu_kernel_params p) |
|
{ |
|
|
|
static_assert(up == 1 || up == 2 || up == 4, "only up=1, up=2, up=4 scales supported"); |
|
static_assert(down == 1 || down == 2 || down == 4, "only down=1, down=2, down=4 scales supported"); |
|
static_assert(fuSize >= up, "upsampling filter size must be at least upsampling factor"); |
|
static_assert(fdSize >= down, "downsampling filter size must be at least downsampling factor"); |
|
static_assert(fuSize % up == 0, "upsampling filter size must be divisible with upsampling factor"); |
|
static_assert(fdSize % down == 0, "downsampling filter size must be divisible with downsampling factor"); |
|
static_assert(fuSize <= MAX_FILTER_SIZE && fdSize <= MAX_FILTER_SIZE, "filter size greater than MAX_FILTER_SIZE"); |
|
static_assert(up != 1 || (fuSize == 1 && (filterMode == MODE_FUFD || filterMode == MODE_FUSD)), "up=1 supported only for 1x1 full filters"); |
|
static_assert(down != 1 || (fdSize == 1 && (filterMode == MODE_FUFD || filterMode == MODE_SUFD)), "down=1 supported only for 1x1 full filters"); |
|
static_assert(!(up == 4 && (filterMode == MODE_FUFD || filterMode == MODE_FUSD)), "full filters not supported for up=4"); |
|
static_assert(!(down == 4 && (filterMode == MODE_FUFD || filterMode == MODE_SUFD)), "full filters not supported for down=4"); |
|
|
|
|
|
typedef typename InternalType<T>::scalar_t scalar_t; |
|
typedef typename InternalType<T>::vec2_t vec2_t; |
|
typedef typename InternalType<T>::vec4_t vec4_t; |
|
const int tileUpW = (tileOutW * down + (fdSize - 1) - (down - 1) + 3) & ~3; |
|
const int tileUpH = tileOutH * down + (fdSize - 1) - (down - 1); |
|
const int tileInW = CEIL_DIV(tileUpW + (fuSize - 1), up); |
|
const int tileInH = CEIL_DIV(tileUpH + (fuSize - 1), up); |
|
const int tileUpH_up = CEIL_DIV(tileUpH, up) * up; |
|
const int tileInH_up = CEIL_DIV(tileUpH_up + (fuSize - 1), up); |
|
|
|
|
|
const bool downInline = (down == 1) && ((up == 1 && filterMode == MODE_FUFD) || (up == 2 && filterMode == MODE_SUFD)); |
|
|
|
|
|
const int szIn = tileInH_up * tileInW; |
|
const int szUpX = tileInH_up * tileUpW; |
|
const int szUpXY = downInline ? 0 : (tileUpH * tileUpW); |
|
const int szDownX = tileUpH * tileOutW; |
|
|
|
|
|
const int s_buf0_size_base = |
|
(filterMode == MODE_SUSD) ? MAX(szIn, szUpXY) : |
|
(filterMode == MODE_FUSD) ? MAX(szIn, szDownX) : |
|
(filterMode == MODE_SUFD) ? MAX(szIn, szUpXY) : |
|
(filterMode == MODE_FUFD) ? szIn : |
|
-1; |
|
const int s_buf1_size_base = |
|
(filterMode == MODE_SUSD) ? MAX(szUpX, szDownX) : |
|
(filterMode == MODE_FUSD) ? szUpXY : |
|
(filterMode == MODE_SUFD) ? szUpX : |
|
(filterMode == MODE_FUFD) ? szUpXY : |
|
-1; |
|
|
|
|
|
const int s_buf0_size = (s_buf0_size_base + 3) & ~3; |
|
const int s_buf1_size = (s_buf1_size_base + 3) & ~3; |
|
|
|
|
|
static_assert((s_buf0_size + s_buf1_size) * sizeof(scalar_t) <= (sharedKB << 10), "shared memory overflow"); |
|
|
|
|
|
scalar_t* s_buf0; |
|
scalar_t* s_buf1; |
|
if (sharedKB <= 48) |
|
{ |
|
|
|
__shared__ scalar_t s_buf0_st[(sharedKB > 48) ? (1<<24) : (s_buf0_size + s_buf1_size)]; |
|
s_buf0 = s_buf0_st; |
|
s_buf1 = s_buf0 + s_buf0_size; |
|
} |
|
else |
|
{ |
|
|
|
s_buf0 = (scalar_t*)s_buf_raw; |
|
s_buf1 = s_buf0 + s_buf0_size; |
|
} |
|
|
|
|
|
scalar_t* s_tileIn; |
|
scalar_t* s_tileUpX; |
|
scalar_t* s_tileUpXY; |
|
scalar_t* s_tileDownX; |
|
if (filterMode == MODE_SUSD) |
|
{ |
|
s_tileIn = s_buf0; |
|
s_tileUpX = s_buf1; |
|
s_tileUpXY = s_buf0; |
|
s_tileDownX = s_buf1; |
|
} |
|
else if (filterMode == MODE_FUSD) |
|
{ |
|
s_tileIn = s_buf0; |
|
s_tileUpXY = s_buf1; |
|
s_tileDownX = s_buf0; |
|
} |
|
else if (filterMode == MODE_SUFD) |
|
{ |
|
s_tileIn = s_buf0; |
|
s_tileUpX = s_buf1; |
|
s_tileUpXY = s_buf0; |
|
} |
|
else if (filterMode == MODE_FUFD) |
|
{ |
|
s_tileIn = s_buf0; |
|
s_tileUpXY = s_buf1; |
|
} |
|
|
|
|
|
int channelIdx = blockIdx.z + p.blockZofs; |
|
int batchIdx = channelIdx / p.yShape.z; |
|
channelIdx -= batchIdx * p.yShape.z; |
|
|
|
|
|
index_t mapOfsOut = channelIdx * get_stride<index_t>(p.yStride.z) + batchIdx * get_stride<index_t>(p.yStride.w); |
|
|
|
|
|
uint32_t signXo = ((threadIdx.x + p.sOfs.x) << 1) & 6; |
|
|
|
|
|
#pragma unroll 1 |
|
for (int tileIdx = 0; !enableXrep || (tileIdx < MIN(p.tilesXrep, p.tilesXdim - p.tilesXrep * blockIdx.y)); tileIdx++) |
|
{ |
|
|
|
int tileX = enableXrep ? blockIdx.y * p.tilesXrep + tileIdx : blockIdx.x; |
|
int tileOutX = tileX * tileOutW; |
|
int tileOutY = (enableXrep ? blockIdx.x : blockIdx.y) * tileOutH; |
|
|
|
|
|
int tmpX = tileOutX * down - p.pad0.x; |
|
int tmpY = tileOutY * down - p.pad0.y; |
|
int tileInX = CEIL_DIV(tmpX, up); |
|
int tileInY = CEIL_DIV(tmpY, up); |
|
const int phaseInX = tileInX * up - tmpX; |
|
const int phaseInY = tileInY * up - tmpY; |
|
|
|
|
|
if (enableXrep && tileIdx > 0 && (filterMode == MODE_FUSD || (filterMode == MODE_SUFD && !downInline) || (filterMode == MODE_FUFD && downInline))) |
|
__syncthreads(); |
|
|
|
|
|
scalar_t b = (scalar_t)*(const T*)((const char*)p.b + (channelIdx * get_stride<index_t>(p.bStride))); |
|
index_t mapOfsIn = channelIdx * get_stride<index_t>(p.xStride.z) + batchIdx * get_stride<index_t>(p.xStride.w); |
|
int idx = threadIdx.x; |
|
const int loopCountIN = CEIL_DIV(tileInW * tileInH, threadsPerBlock); |
|
#pragma unroll |
|
for (int loop = 0; loop < loopCountIN; loop++) |
|
{ |
|
int relInX, relInY; |
|
fast_div_mod<tileInW>(relInX, relInY, idx); |
|
int inX = tileInX + relInX; |
|
int inY = tileInY + relInY; |
|
scalar_t v = 0; |
|
|
|
if ((uint32_t)inX < p.xShape.x && (uint32_t)inY < p.xShape.y) |
|
v = (scalar_t)*((const T*)((const char*)p.x + (inX * get_stride<index_t>(p.xStride.x) + inY * get_stride<index_t>(p.xStride.y) + mapOfsIn))) + b; |
|
|
|
bool skip = (loop == loopCountIN-1) && (idx >= tileInW * tileInH); |
|
if (!skip) |
|
s_tileIn[idx] = v; |
|
|
|
idx += threadsPerBlock; |
|
} |
|
|
|
if (filterMode == MODE_SUSD || filterMode == MODE_SUFD) |
|
{ |
|
|
|
__syncthreads(); |
|
if (up == 4) |
|
{ |
|
for (int idx = threadIdx.x*up; idx < tileUpW * tileInH; idx += blockDim.x*up) |
|
{ |
|
int relUpX0, relInY; |
|
fast_div_mod<tileUpW>(relUpX0, relInY, idx); |
|
int relInX0 = relUpX0 / up; |
|
int src0 = relInX0 + tileInW * relInY; |
|
int dst = relInY * tileUpW + relUpX0; |
|
vec4_t v = InternalType<T>::zero_vec4(); |
|
scalar_t a = s_tileIn[src0]; |
|
if (phaseInX == 0) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileIn[src0 + step + 1]; |
|
v.y += a * (scalar_t)c_fu[step * up + 3]; |
|
v.z += a * (scalar_t)c_fu[step * up + 2]; |
|
v.w += a * (scalar_t)c_fu[step * up + 1]; |
|
} |
|
} |
|
else if (phaseInX == 1) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 1]; |
|
v.y += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileIn[src0 + step + 1]; |
|
v.z += a * (scalar_t)c_fu[step * up + 3]; |
|
v.w += a * (scalar_t)c_fu[step * up + 2]; |
|
} |
|
} |
|
else if (phaseInX == 2) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 2]; |
|
v.y += a * (scalar_t)c_fu[step * up + 1]; |
|
v.z += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileIn[src0 + step + 1]; |
|
v.w += a * (scalar_t)c_fu[step * up + 3]; |
|
} |
|
} |
|
else |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 3]; |
|
v.y += a * (scalar_t)c_fu[step * up + 2]; |
|
v.z += a * (scalar_t)c_fu[step * up + 1]; |
|
v.w += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileIn[src0 + step + 1]; |
|
} |
|
} |
|
s_tileUpX[dst+0] = v.x; |
|
s_tileUpX[dst+1] = v.y; |
|
s_tileUpX[dst+2] = v.z; |
|
s_tileUpX[dst+3] = v.w; |
|
} |
|
} |
|
else if (up == 2) |
|
{ |
|
bool p0 = (phaseInX == 0); |
|
for (int idx = threadIdx.x*up; idx < tileUpW * tileInH; idx += blockDim.x*up) |
|
{ |
|
int relUpX0, relInY; |
|
fast_div_mod<tileUpW>(relUpX0, relInY, idx); |
|
int relInX0 = relUpX0 / up; |
|
int src0 = relInX0 + tileInW * relInY; |
|
int dst = relInY * tileUpW + relUpX0; |
|
vec2_t v = InternalType<T>::zero_vec2(); |
|
scalar_t a = s_tileIn[src0]; |
|
if (p0) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileIn[src0 + step + 1]; |
|
v.y += a * (scalar_t)c_fu[step * up + 1]; |
|
} |
|
} |
|
else |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 1]; |
|
v.y += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileIn[src0 + step + 1]; |
|
} |
|
} |
|
s_tileUpX[dst+0] = v.x; |
|
s_tileUpX[dst+1] = v.y; |
|
} |
|
} |
|
|
|
|
|
|
|
__syncthreads(); |
|
int groupMask = 15 << ((threadIdx.x & 31) & ~3); |
|
int minY = tileOutY ? (tileOutY - tileOutH) * down + tileUpH : 0; |
|
int sShapeMaxY = MIN(p.sShape.y, tileOutY * down + tileUpH); |
|
if (up == 4) |
|
{ |
|
minY -= 3; |
|
for (int idx = threadIdx.x; idx < tileUpW * tileUpH_up / up; idx += blockDim.x) |
|
{ |
|
int relUpX, relInY0; |
|
fast_div_mod<tileUpW>(relUpX, relInY0, idx); |
|
int relUpY0 = relInY0 * up; |
|
int src0 = relInY0 * tileUpW + relUpX; |
|
int dst = relUpY0 * tileUpW + relUpX; |
|
vec4_t v = InternalType<T>::zero_vec4(); |
|
|
|
scalar_t a = s_tileUpX[src0]; |
|
if (phaseInY == 0) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileUpX[src0 + (step + 1) * tileUpW]; |
|
v.y += a * (scalar_t)c_fu[step * up + 3]; |
|
v.z += a * (scalar_t)c_fu[step * up + 2]; |
|
v.w += a * (scalar_t)c_fu[step * up + 1]; |
|
} |
|
} |
|
else if (phaseInY == 1) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 1]; |
|
v.y += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileUpX[src0 + (step + 1) * tileUpW]; |
|
v.z += a * (scalar_t)c_fu[step * up + 3]; |
|
v.w += a * (scalar_t)c_fu[step * up + 2]; |
|
} |
|
} |
|
else if (phaseInY == 2) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 2]; |
|
v.y += a * (scalar_t)c_fu[step * up + 1]; |
|
v.z += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileUpX[src0 + (step + 1) * tileUpW]; |
|
v.w += a * (scalar_t)c_fu[step * up + 3]; |
|
} |
|
} |
|
else |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 3]; |
|
v.y += a * (scalar_t)c_fu[step * up + 2]; |
|
v.z += a * (scalar_t)c_fu[step * up + 1]; |
|
v.w += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileUpX[src0 + (step + 1) * tileUpW]; |
|
} |
|
} |
|
|
|
int x = tileOutX * down + relUpX; |
|
int y = tileOutY * down + relUpY0; |
|
int signX = x + p.sOfs.x; |
|
int signY = y + p.sOfs.y; |
|
int signZ = blockIdx.z + p.blockZofs; |
|
int signXb = signX >> 2; |
|
index_t si0 = signXb + p.sShape.x * (signY + (index_t)p.sShape.y * signZ); |
|
index_t si1 = si0 + p.sShape.x; |
|
index_t si2 = si0 + p.sShape.x * 2; |
|
index_t si3 = si0 + p.sShape.x * 3; |
|
|
|
v.x *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.y *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.z *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.w *= (scalar_t)((float)up * (float)up * p.gain); |
|
|
|
if (signWrite) |
|
{ |
|
if (!enableWriteSkip) |
|
{ |
|
|
|
int sx = __float_as_uint(v.x) >> 31 << 0; |
|
int sy = __float_as_uint(v.y) >> 31 << 8; |
|
int sz = __float_as_uint(v.z) >> 31 << 16; |
|
int sw = __float_as_uint(v.w) >> 31 << 24; |
|
if (sx) v.x *= p.slope; |
|
if (sy) v.y *= p.slope; |
|
if (sz) v.z *= p.slope; |
|
if (sw) v.w *= p.slope; |
|
if (fabsf(v.x) > p.clamp) { sx = 2 << 0; v.x = InternalType<T>::clamp(v.x, p.clamp); } |
|
if (fabsf(v.y) > p.clamp) { sy = 2 << 8; v.y = InternalType<T>::clamp(v.y, p.clamp); } |
|
if (fabsf(v.z) > p.clamp) { sz = 2 << 16; v.z = InternalType<T>::clamp(v.z, p.clamp); } |
|
if (fabsf(v.w) > p.clamp) { sw = 2 << 24; v.w = InternalType<T>::clamp(v.w, p.clamp); } |
|
|
|
if ((uint32_t)signXb < p.swLimit && signY >= minY) |
|
{ |
|
|
|
uint32_t s = sx + sy + sw + sz; |
|
s <<= (signX & 3) << 1; |
|
s |= __shfl_xor_sync(groupMask, s, 1); |
|
s |= __shfl_xor_sync(groupMask, s, 2); |
|
|
|
|
|
if ((uint32_t)(signY + 0) < sShapeMaxY) { p.s[si0] = (unsigned char)(s >> 0); } |
|
if ((uint32_t)(signY + 1) < sShapeMaxY) { p.s[si1] = (unsigned char)(s >> 8); } |
|
if ((uint32_t)(signY + 2) < sShapeMaxY) { p.s[si2] = (unsigned char)(s >> 16); } |
|
if ((uint32_t)(signY + 3) < sShapeMaxY) { p.s[si3] = (unsigned char)(s >> 24); } |
|
} |
|
} |
|
else |
|
{ |
|
|
|
if ((uint32_t)signXb < p.swLimit && signY >= minY) |
|
{ |
|
int sx = __float_as_uint(v.x) >> 31 << 0; |
|
int sy = __float_as_uint(v.y) >> 31 << 8; |
|
int sz = __float_as_uint(v.z) >> 31 << 16; |
|
int sw = __float_as_uint(v.w) >> 31 << 24; |
|
if (sx) v.x *= p.slope; |
|
if (sy) v.y *= p.slope; |
|
if (sz) v.z *= p.slope; |
|
if (sw) v.w *= p.slope; |
|
if (fabsf(v.x) > p.clamp) { sx = 2 << 0; v.x = InternalType<T>::clamp(v.x, p.clamp); } |
|
if (fabsf(v.y) > p.clamp) { sy = 2 << 8; v.y = InternalType<T>::clamp(v.y, p.clamp); } |
|
if (fabsf(v.z) > p.clamp) { sz = 2 << 16; v.z = InternalType<T>::clamp(v.z, p.clamp); } |
|
if (fabsf(v.w) > p.clamp) { sw = 2 << 24; v.w = InternalType<T>::clamp(v.w, p.clamp); } |
|
|
|
|
|
uint32_t s = sx + sy + sw + sz; |
|
s <<= (signX & 3) << 1; |
|
s |= __shfl_xor_sync(groupMask, s, 1); |
|
s |= __shfl_xor_sync(groupMask, s, 2); |
|
|
|
|
|
if ((uint32_t)(signY + 0) < sShapeMaxY) { p.s[si0] = (unsigned char)(s >> 0); } |
|
if ((uint32_t)(signY + 1) < sShapeMaxY) { p.s[si1] = (unsigned char)(s >> 8); } |
|
if ((uint32_t)(signY + 2) < sShapeMaxY) { p.s[si2] = (unsigned char)(s >> 16); } |
|
if ((uint32_t)(signY + 3) < sShapeMaxY) { p.s[si3] = (unsigned char)(s >> 24); } |
|
} |
|
else |
|
{ |
|
|
|
if (v.x < 0.f) v.x *= p.slope; v.x = InternalType<T>::clamp(v.x, p.clamp); |
|
if (v.y < 0.f) v.y *= p.slope; v.y = InternalType<T>::clamp(v.y, p.clamp); |
|
if (v.z < 0.f) v.z *= p.slope; v.z = InternalType<T>::clamp(v.z, p.clamp); |
|
if (v.w < 0.f) v.w *= p.slope; v.w = InternalType<T>::clamp(v.w, p.clamp); |
|
} |
|
} |
|
} |
|
else if (signRead) |
|
{ |
|
if ((uint32_t)signXb < p.swLimit) |
|
{ |
|
int ss = (signX & 3) << 1; |
|
if ((uint32_t)(signY + 0) < p.sShape.y) { int s = p.s[si0] >> ss; if (s & 1) v.x *= p.slope; if (s & 2) v.x = 0.f; } |
|
if ((uint32_t)(signY + 1) < p.sShape.y) { int s = p.s[si1] >> ss; if (s & 1) v.y *= p.slope; if (s & 2) v.y = 0.f; } |
|
if ((uint32_t)(signY + 2) < p.sShape.y) { int s = p.s[si2] >> ss; if (s & 1) v.z *= p.slope; if (s & 2) v.z = 0.f; } |
|
if ((uint32_t)(signY + 3) < p.sShape.y) { int s = p.s[si3] >> ss; if (s & 1) v.w *= p.slope; if (s & 2) v.w = 0.f; } |
|
} |
|
} |
|
else |
|
{ |
|
if (v.x < 0.f) v.x *= p.slope; v.x = InternalType<T>::clamp(v.x, p.clamp); |
|
if (v.y < 0.f) v.y *= p.slope; v.y = InternalType<T>::clamp(v.y, p.clamp); |
|
if (v.z < 0.f) v.z *= p.slope; v.z = InternalType<T>::clamp(v.z, p.clamp); |
|
if (v.w < 0.f) v.w *= p.slope; v.w = InternalType<T>::clamp(v.w, p.clamp); |
|
} |
|
|
|
s_tileUpXY[dst + 0 * tileUpW] = v.x; |
|
if (relUpY0 + 1 < tileUpH) s_tileUpXY[dst + 1 * tileUpW] = v.y; |
|
if (relUpY0 + 2 < tileUpH) s_tileUpXY[dst + 2 * tileUpW] = v.z; |
|
if (relUpY0 + 3 < tileUpH) s_tileUpXY[dst + 3 * tileUpW] = v.w; |
|
} |
|
} |
|
else if (up == 2) |
|
{ |
|
minY -= 1; |
|
for (int idx = threadIdx.x; idx < tileUpW * tileUpH_up / up; idx += blockDim.x) |
|
{ |
|
int relUpX, relInY0; |
|
fast_div_mod<tileUpW>(relUpX, relInY0, idx); |
|
int relUpY0 = relInY0 * up; |
|
int src0 = relInY0 * tileUpW + relUpX; |
|
int dst = relUpY0 * tileUpW + relUpX; |
|
vec2_t v = InternalType<T>::zero_vec2(); |
|
|
|
scalar_t a = s_tileUpX[src0]; |
|
if (phaseInY == 0) |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileUpX[src0 + (step + 1) * tileUpW]; |
|
v.y += a * (scalar_t)c_fu[step * up + 1]; |
|
} |
|
} |
|
else |
|
{ |
|
#pragma unroll |
|
for (int step = 0; step < fuSize / up; step++) |
|
{ |
|
v.x += a * (scalar_t)c_fu[step * up + 1]; |
|
v.y += a * (scalar_t)c_fu[step * up + 0]; |
|
a = s_tileUpX[src0 + (step + 1) * tileUpW]; |
|
} |
|
} |
|
|
|
int x = tileOutX * down + relUpX; |
|
int y = tileOutY * down + relUpY0; |
|
int signX = x + p.sOfs.x; |
|
int signY = y + p.sOfs.y; |
|
int signZ = blockIdx.z + p.blockZofs; |
|
int signXb = signX >> 2; |
|
index_t si0 = signXb + p.sShape.x * (signY + (index_t)p.sShape.y * signZ); |
|
index_t si1 = si0 + p.sShape.x; |
|
|
|
v.x *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.y *= (scalar_t)((float)up * (float)up * p.gain); |
|
|
|
if (signWrite) |
|
{ |
|
if (!enableWriteSkip) |
|
{ |
|
|
|
int sx = __float_as_uint(v.x) >> 31 << 0; |
|
int sy = __float_as_uint(v.y) >> 31 << 8; |
|
if (sx) v.x *= p.slope; |
|
if (sy) v.y *= p.slope; |
|
if (fabsf(v.x) > p.clamp) { sx = 2 << 0; v.x = InternalType<T>::clamp(v.x, p.clamp); } |
|
if (fabsf(v.y) > p.clamp) { sy = 2 << 8; v.y = InternalType<T>::clamp(v.y, p.clamp); } |
|
|
|
if ((uint32_t)signXb < p.swLimit && signY >= minY) |
|
{ |
|
|
|
int s = sx + sy; |
|
s <<= signXo; |
|
s |= __shfl_xor_sync(groupMask, s, 1); |
|
s |= __shfl_xor_sync(groupMask, s, 2); |
|
|
|
|
|
if ((uint32_t)(signY + 0) < sShapeMaxY) { p.s[si0] = (unsigned char)(s >> 0); } |
|
if ((uint32_t)(signY + 1) < sShapeMaxY) { p.s[si1] = (unsigned char)(s >> 8); } |
|
} |
|
} |
|
else |
|
{ |
|
|
|
if ((uint32_t)signXb < p.swLimit && signY >= minY) |
|
{ |
|
int sx = __float_as_uint(v.x) >> 31 << 0; |
|
int sy = __float_as_uint(v.y) >> 31 << 8; |
|
if (sx) v.x *= p.slope; |
|
if (sy) v.y *= p.slope; |
|
if (fabsf(v.x) > p.clamp) { sx = 2 << 0; v.x = InternalType<T>::clamp(v.x, p.clamp); } |
|
if (fabsf(v.y) > p.clamp) { sy = 2 << 8; v.y = InternalType<T>::clamp(v.y, p.clamp); } |
|
|
|
|
|
int s = sx + sy; |
|
s <<= signXo; |
|
s |= __shfl_xor_sync(groupMask, s, 1); |
|
s |= __shfl_xor_sync(groupMask, s, 2); |
|
|
|
|
|
if ((uint32_t)(signY + 0) < sShapeMaxY) { p.s[si0] = (unsigned char)(s >> 0); } |
|
if ((uint32_t)(signY + 1) < sShapeMaxY) { p.s[si1] = (unsigned char)(s >> 8); } |
|
} |
|
else |
|
{ |
|
|
|
if (v.x < 0.f) v.x *= p.slope; v.x = InternalType<T>::clamp(v.x, p.clamp); |
|
if (v.y < 0.f) v.y *= p.slope; v.y = InternalType<T>::clamp(v.y, p.clamp); |
|
} |
|
} |
|
} |
|
else if (signRead) |
|
{ |
|
if ((uint32_t)signXb < p.swLimit) |
|
{ |
|
if ((uint32_t)(signY + 0) < p.sShape.y) { int s = p.s[si0] >> signXo; if (s & 1) v.x *= p.slope; if (s & 2) v.x = 0.f; } |
|
if ((uint32_t)(signY + 1) < p.sShape.y) { int s = p.s[si1] >> signXo; if (s & 1) v.y *= p.slope; if (s & 2) v.y = 0.f; } |
|
} |
|
} |
|
else |
|
{ |
|
if (v.x < 0.f) v.x *= p.slope; v.x = InternalType<T>::clamp(v.x, p.clamp); |
|
if (v.y < 0.f) v.y *= p.slope; v.y = InternalType<T>::clamp(v.y, p.clamp); |
|
} |
|
|
|
if (!downInline) |
|
{ |
|
|
|
s_tileUpXY[dst] = v.x; |
|
if (relUpY0 < tileUpH - 1) |
|
s_tileUpXY[dst + tileUpW] = v.y; |
|
} |
|
else |
|
{ |
|
|
|
if ((uint32_t)x < p.yShape.x) |
|
{ |
|
int ymax = MIN(p.yShape.y, tileUpH + tileOutY * down); |
|
index_t ofs = x * get_stride<index_t>(p.yStride.x) + y * get_stride<index_t>(p.yStride.y) + mapOfsOut; |
|
if ((uint32_t)y + 0 < p.yShape.y) *((T*)((char*)p.y + ofs)) = (T)(v.x * (scalar_t)c_fd[0]); |
|
if ((uint32_t)y + 1 < ymax) *((T*)((char*)p.y + ofs + get_stride<index_t>(p.yStride.y))) = (T)(v.y * (scalar_t)c_fd[0]); |
|
} |
|
} |
|
} |
|
} |
|
} |
|
else if (filterMode == MODE_FUSD || filterMode == MODE_FUFD) |
|
{ |
|
|
|
|
|
if (up == 2) |
|
{ |
|
|
|
__syncthreads(); |
|
int minY = tileOutY ? (tileOutY - tileOutH) * down + tileUpH + p.sOfs.y : 0; |
|
for (int idx = threadIdx.x * 4; idx < tileUpW * tileUpH; idx += blockDim.x * 4) |
|
{ |
|
int relUpX0, relUpY0; |
|
fast_div_mod<tileUpW>(relUpX0, relUpY0, idx); |
|
int relInX0 = CEIL_DIV(relUpX0 - phaseInX, up); |
|
int relInY0 = CEIL_DIV(relUpY0 - phaseInY, up); |
|
int src0 = relInX0 + tileInW * relInY0; |
|
int tap0y = (relInY0 * up + phaseInY - relUpY0); |
|
|
|
#define X_LOOP(TAPY, PX) \ |
|
for (int sx = 0; sx < fuSize / up; sx++) \ |
|
{ \ |
|
v.x += a * (scalar_t)c_fu[(sx * up + (((PX) - 0) & (up - 1))) + (sy * up + (TAPY)) * MAX_FILTER_SIZE]; \ |
|
v.z += b * (scalar_t)c_fu[(sx * up + (((PX) - 0) & (up - 1))) + (sy * up + (TAPY)) * MAX_FILTER_SIZE]; if ((PX) == 0) { a = b; b = s_tileIn[src0 + 2 + sx + sy * tileInW]; } \ |
|
v.y += a * (scalar_t)c_fu[(sx * up + (((PX) - 1) & (up - 1))) + (sy * up + (TAPY)) * MAX_FILTER_SIZE]; \ |
|
v.w += b * (scalar_t)c_fu[(sx * up + (((PX) - 1) & (up - 1))) + (sy * up + (TAPY)) * MAX_FILTER_SIZE]; if ((PX) == 1) { a = b; b = s_tileIn[src0 + 2 + sx + sy * tileInW]; } \ |
|
} |
|
|
|
vec4_t v = InternalType<T>::zero_vec4(); |
|
if (tap0y == 0 && phaseInX == 0) |
|
#pragma unroll |
|
for (int sy = 0; sy < fuSize / up; sy++) { scalar_t a = s_tileIn[src0 + sy * tileInW]; scalar_t b = s_tileIn[src0 + sy * tileInW + 1]; |
|
#pragma unroll |
|
X_LOOP(0, 0) } |
|
if (tap0y == 0 && phaseInX == 1) |
|
#pragma unroll |
|
for (int sy = 0; sy < fuSize / up; sy++) { scalar_t a = s_tileIn[src0 + sy * tileInW]; scalar_t b = s_tileIn[src0 + sy * tileInW + 1]; |
|
#pragma unroll |
|
X_LOOP(0, 1) } |
|
if (tap0y == 1 && phaseInX == 0) |
|
#pragma unroll |
|
for (int sy = 0; sy < fuSize / up; sy++) { scalar_t a = s_tileIn[src0 + sy * tileInW]; scalar_t b = s_tileIn[src0 + sy * tileInW + 1]; |
|
#pragma unroll |
|
X_LOOP(1, 0) } |
|
if (tap0y == 1 && phaseInX == 1) |
|
#pragma unroll |
|
for (int sy = 0; sy < fuSize / up; sy++) { scalar_t a = s_tileIn[src0 + sy * tileInW]; scalar_t b = s_tileIn[src0 + sy * tileInW + 1]; |
|
#pragma unroll |
|
X_LOOP(1, 1) } |
|
|
|
#undef X_LOOP |
|
|
|
int x = tileOutX * down + relUpX0; |
|
int y = tileOutY * down + relUpY0; |
|
int signX = x + p.sOfs.x; |
|
int signY = y + p.sOfs.y; |
|
int signZ = blockIdx.z + p.blockZofs; |
|
int signXb = signX >> 2; |
|
index_t si = signXb + p.sShape.x * (signY + (index_t)p.sShape.y * signZ); |
|
|
|
v.x *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.y *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.z *= (scalar_t)((float)up * (float)up * p.gain); |
|
v.w *= (scalar_t)((float)up * (float)up * p.gain); |
|
|
|
if (signWrite) |
|
{ |
|
if (!enableWriteSkip) |
|
{ |
|
|
|
int sx = __float_as_uint(v.x) >> 31; |
|
int sy = __float_as_uint(v.y) >> 31; |
|
int sz = __float_as_uint(v.z) >> 31; |
|
int sw = __float_as_uint(v.w) >> 31; |
|
if (sx) v.x *= p.slope; if (fabsf(v.x) > p.clamp) { sx = 2; v.x = InternalType<T>::clamp(v.x, p.clamp); } |
|
if (sy) v.y *= p.slope; if (fabsf(v.y) > p.clamp) { sy = 2; v.y = InternalType<T>::clamp(v.y, p.clamp); } |
|
if (sz) v.z *= p.slope; if (fabsf(v.z) > p.clamp) { sz = 2; v.z = InternalType<T>::clamp(v.z, p.clamp); } |
|
if (sw) v.w *= p.slope; if (fabsf(v.w) > p.clamp) { sw = 2; v.w = InternalType<T>::clamp(v.w, p.clamp); } |
|
|
|
if ((uint32_t)signXb < p.swLimit && (uint32_t)signY < p.sShape.y && signY >= minY) |
|
{ |
|
p.s[si] = sx + (sy << 2) + (sz << 4) + (sw << 6); |
|
} |
|
} |
|
else |
|
{ |
|
|
|
if ((uint32_t)signXb < p.swLimit && (uint32_t)signY < p.sShape.y && signY >= minY) |
|
{ |
|
int sx = __float_as_uint(v.x) >> 31; |
|
int sy = __float_as_uint(v.y) >> 31; |
|
int sz = __float_as_uint(v.z) >> 31; |
|
int sw = __float_as_uint(v.w) >> 31; |
|
if (sx) v.x *= p.slope; if (fabsf(v.x) > p.clamp) { sx = 2; v.x = InternalType<T>::clamp(v.x, p.clamp); } |
|
if (sy) v.y *= p.slope; if (fabsf(v.y) > p.clamp) { sy = 2; v.y = InternalType<T>::clamp(v.y, p.clamp); } |
|
if (sz) v.z *= p.slope; if (fabsf(v.z) > p.clamp) { sz = 2; v.z = InternalType<T>::clamp(v.z, p.clamp); } |
|
if (sw) v.w *= p.slope; if (fabsf(v.w) > p.clamp) { sw = 2; v.w = InternalType<T>::clamp(v.w, p.clamp); } |
|
|
|
p.s[si] = sx + (sy << 2) + (sz << 4) + (sw << 6); |
|
} |
|
else |
|
{ |
|
|
|
if (v.x < 0.f) v.x *= p.slope; v.x = InternalType<T>::clamp(v.x, p.clamp); |
|
if (v.y < 0.f) v.y *= p.slope; v.y = InternalType<T>::clamp(v.y, p.clamp); |
|
if (v.z < 0.f) v.z *= p.slope; v.z = InternalType<T>::clamp(v.z, p.clamp); |
|
if (v.w < 0.f) v.w *= p.slope; v.w = InternalType<T>::clamp(v.w, p.clamp); |
|
} |
|
} |
|
} |
|
else if (signRead) |
|
{ |
|
if ((uint32_t)signY < p.sShape.y) |
|
{ |
|
int s = 0; |
|
if ((uint32_t)signXb < p.swLimit) s = p.s[si]; |
|
if ((uint32_t)signXb + 1 < p.swLimit) s |= p.s[si + 1] << 8; |
|
s >>= (signX & 3) << 1; |
|
if (s & 0x01) v.x *= p.slope; if (s & 0x02) v.x = 0.f; |
|
if (s & 0x04) v.y *= p.slope; if (s & 0x08) v.y = 0.f; |
|
if (s & 0x10) v.z *= p.slope; if (s & 0x20) v.z = 0.f; |
|
if (s & 0x40) v.w *= p.slope; if (s & 0x80) v.w = 0.f; |
|
} |
|
} |
|
else |
|
{ |
|
if (v.x < 0.f) v.x *= p.slope; v.x = InternalType<T>::clamp(v.x, p.clamp); |
|
if (v.y < 0.f) v.y *= p.slope; v.y = InternalType<T>::clamp(v.y, p.clamp); |
|
if (v.z < 0.f) v.z *= p.slope; v.z = InternalType<T>::clamp(v.z, p.clamp); |
|
if (v.w < 0.f) v.w *= p.slope; v.w = InternalType<T>::clamp(v.w, p.clamp); |
|
} |
|
|
|
s_tileUpXY[idx + 0] = v.x; |
|
s_tileUpXY[idx + 1] = v.y; |
|
s_tileUpXY[idx + 2] = v.z; |
|
s_tileUpXY[idx + 3] = v.w; |
|
} |
|
} |
|
else if (up == 1) |
|
{ |
|
__syncthreads(); |
|
uint32_t groupMask = 15 << ((threadIdx.x & 31) & ~3); |
|
int minY = tileOutY ? (tileOutY - tileOutH) * down + tileUpH : 0; |
|
for (int idx = threadIdx.x; idx < tileUpW * tileUpH; idx += blockDim.x) |
|
{ |
|
int relUpX0, relUpY0; |
|
fast_div_mod<tileUpW>(relUpX0, relUpY0, idx); |
|
scalar_t v = s_tileIn[idx] * (scalar_t)c_fu[0]; |
|
|
|
int x = tileOutX * down + relUpX0; |
|
int y = tileOutY * down + relUpY0; |
|
int signX = x + p.sOfs.x; |
|
int signY = y + p.sOfs.y; |
|
int signZ = blockIdx.z + p.blockZofs; |
|
int signXb = signX >> 2; |
|
index_t si = signXb + p.sShape.x * (signY + (index_t)p.sShape.y * signZ); |
|
v *= (scalar_t)((float)up * (float)up * p.gain); |
|
|
|
if (signWrite) |
|
{ |
|
if (!enableWriteSkip) |
|
{ |
|
|
|
uint32_t s = 0; |
|
uint32_t signXbit = (1u << signXo); |
|
if (v < 0.f) |
|
{ |
|
s = signXbit; |
|
v *= p.slope; |
|
} |
|
if (fabsf(v) > p.clamp) |
|
{ |
|
s = signXbit * 2; |
|
v = InternalType<T>::clamp(v, p.clamp); |
|
} |
|
if ((uint32_t)signXb < p.swLimit && (uint32_t)signY < p.sShape.y && signY >= minY) |
|
{ |
|
s += __shfl_xor_sync(groupMask, s, 1); |
|
s += __shfl_xor_sync(groupMask, s, 2); |
|
p.s[si] = s; |
|
} |
|
} |
|
else |
|
{ |
|
|
|
if ((uint32_t)signXb < p.swLimit && (uint32_t)signY < p.sShape.y && signY >= minY) |
|
{ |
|
uint32_t s = 0; |
|
uint32_t signXbit = (1u << signXo); |
|
if (v < 0.f) |
|
{ |
|
s = signXbit; |
|
v *= p.slope; |
|
} |
|
if (fabsf(v) > p.clamp) |
|
{ |
|
s = signXbit * 2; |
|
v = InternalType<T>::clamp(v, p.clamp); |
|
} |
|
s += __shfl_xor_sync(groupMask, s, 1); |
|
s += __shfl_xor_sync(groupMask, s, 2); |
|
p.s[si] = s; |
|
} |
|
else |
|
{ |
|
|
|
if (v < 0.f) v *= p.slope; |
|
v = InternalType<T>::clamp(v, p.clamp); |
|
} |
|
} |
|
} |
|
else if (signRead) |
|
{ |
|
|
|
if ((uint32_t)signXb < p.swLimit && (uint32_t)signY < p.sShape.y) |
|
{ |
|
int s = p.s[si]; |
|
s >>= signXo; |
|
if (s & 1) v *= p.slope; |
|
if (s & 2) v = 0.f; |
|
} |
|
} |
|
else |
|
{ |
|
if (v < 0.f) v *= p.slope; |
|
v = InternalType<T>::clamp(v, p.clamp); |
|
} |
|
|
|
if (!downInline) |
|
s_tileUpXY[idx] = v; |
|
else if ((uint32_t)x < p.yShape.x && (uint32_t)y < p.yShape.y) |
|
*((T*)((char*)p.y + (x * get_stride<index_t>(p.yStride.x) + y * get_stride<index_t>(p.yStride.y) + mapOfsOut))) = (T)(v * (scalar_t)c_fd[0]); |
|
} |
|
} |
|
} |
|
|
|
|
|
if (filterMode == MODE_SUSD || filterMode == MODE_FUSD) |
|
{ |
|
|
|
__syncthreads(); |
|
if (down == 4 && tileOutW % 4 == 0) |
|
{ |
|
|
|
for (int idx = threadIdx.x * 4; idx < tileOutW * tileUpH; idx += blockDim.x * 4) |
|
{ |
|
int relOutX0, relUpY; |
|
fast_div_mod<tileOutW>(relOutX0, relUpY, idx); |
|
int relUpX0 = relOutX0 * down; |
|
int src0 = relUpY * tileUpW + relUpX0; |
|
vec4_t v = InternalType<T>::zero_vec4(); |
|
#pragma unroll |
|
for (int step = 0; step < fdSize; step++) |
|
{ |
|
v.x += s_tileUpXY[src0 + 0 + step] * (scalar_t)c_fd[step]; |
|
v.y += s_tileUpXY[src0 + 4 + step] * (scalar_t)c_fd[step]; |
|
v.z += s_tileUpXY[src0 + 8 + step] * (scalar_t)c_fd[step]; |
|
v.w += s_tileUpXY[src0 + 12 + step] * (scalar_t)c_fd[step]; |
|
} |
|
s_tileDownX[idx+0] = v.x; |
|
s_tileDownX[idx+1] = v.y; |
|
s_tileDownX[idx+2] = v.z; |
|
s_tileDownX[idx+3] = v.w; |
|
} |
|
} |
|
else if ((down == 2 || down == 4) && (tileOutW % 2 == 0)) |
|
{ |
|
|
|
for (int idx = threadIdx.x * 2; idx < tileOutW * tileUpH; idx += blockDim.x * 2) |
|
{ |
|
int relOutX0, relUpY; |
|
fast_div_mod<tileOutW>(relOutX0, relUpY, idx); |
|
int relUpX0 = relOutX0 * down; |
|
int src0 = relUpY * tileUpW + relUpX0; |
|
vec2_t v = InternalType<T>::zero_vec2(); |
|
#pragma unroll |
|
for (int step = 0; step < fdSize; step++) |
|
{ |
|
v.x += s_tileUpXY[src0 + 0 + step] * (scalar_t)c_fd[step]; |
|
v.y += s_tileUpXY[src0 + down + step] * (scalar_t)c_fd[step]; |
|
} |
|
s_tileDownX[idx+0] = v.x; |
|
s_tileDownX[idx+1] = v.y; |
|
} |
|
} |
|
else |
|
{ |
|
|
|
for (int idx = threadIdx.x; idx < tileOutW * tileUpH; idx += blockDim.x) |
|
{ |
|
int relOutX0, relUpY; |
|
fast_div_mod<tileOutW>(relOutX0, relUpY, idx); |
|
int relUpX0 = relOutX0 * down; |
|
int src = relUpY * tileUpW + relUpX0; |
|
scalar_t v = 0.f; |
|
#pragma unroll |
|
for (int step = 0; step < fdSize; step++) |
|
v += s_tileUpXY[src + step] * (scalar_t)c_fd[step]; |
|
s_tileDownX[idx] = v; |
|
} |
|
} |
|
|
|
|
|
__syncthreads(); |
|
for (int idx = threadIdx.x; idx < tileOutW * tileOutH; idx += blockDim.x) |
|
{ |
|
int relOutX, relOutY0; |
|
fast_div_mod<tileOutW>(relOutX, relOutY0, idx); |
|
int relUpY0 = relOutY0 * down; |
|
int src0 = relUpY0 * tileOutW + relOutX; |
|
scalar_t v = 0; |
|
#pragma unroll |
|
for (int step = 0; step < fdSize; step++) |
|
v += s_tileDownX[src0 + step * tileOutW] * (scalar_t)c_fd[step]; |
|
|
|
int outX = tileOutX + relOutX; |
|
int outY = tileOutY + relOutY0; |
|
|
|
if (outX < p.yShape.x & outY < p.yShape.y) |
|
*((T*)((char*)p.y + (outX * get_stride<index_t>(p.yStride.x) + outY * get_stride<index_t>(p.yStride.y) + mapOfsOut))) = (T)v; |
|
} |
|
} |
|
else if (filterMode == MODE_SUFD || filterMode == MODE_FUFD) |
|
{ |
|
|
|
if (down == 2) |
|
{ |
|
|
|
__syncthreads(); |
|
for (int idx = threadIdx.x * 2; idx < tileOutW * tileOutH; idx += blockDim.x * 2) |
|
{ |
|
int relOutX0, relOutY0; |
|
fast_div_mod<tileOutW>(relOutX0, relOutY0, idx); |
|
int relUpX0 = relOutX0 * down; |
|
int relUpY0 = relOutY0 * down; |
|
int src0 = relUpY0 * tileUpW + relUpX0; |
|
vec2_t v = InternalType<T>::zero_vec2(); |
|
#pragma unroll |
|
for (int sy = 0; sy < fdSize; sy++) |
|
#pragma unroll |
|
for (int sx = 0; sx < fdSize; sx++) |
|
{ |
|
v.x += s_tileUpXY[src0 + 0 + sx + sy * tileUpW] * (scalar_t)c_fd[sx + sy * MAX_FILTER_SIZE]; |
|
v.y += s_tileUpXY[src0 + 2 + sx + sy * tileUpW] * (scalar_t)c_fd[sx + sy * MAX_FILTER_SIZE]; |
|
} |
|
|
|
int outX = tileOutX + relOutX0; |
|
int outY = tileOutY + relOutY0; |
|
if ((uint32_t)outY < p.yShape.y) |
|
{ |
|
index_t ofs = outX * get_stride<index_t>(p.yStride.x) + outY * get_stride<index_t>(p.yStride.y) + mapOfsOut; |
|
if (outX + 0 < p.yShape.x) *((T*)((char*)p.y + ofs)) = (T)v.x; |
|
if (outX + 1 < p.yShape.x) *((T*)((char*)p.y + ofs + get_stride<index_t>(p.yStride.x))) = (T)v.y; |
|
} |
|
} |
|
} |
|
else if (down == 1 && !downInline) |
|
{ |
|
|
|
__syncthreads(); |
|
for (int idx = threadIdx.x; idx < tileOutW * tileOutH; idx += blockDim.x) |
|
{ |
|
int relOutX0, relOutY0; |
|
fast_div_mod<tileOutW>(relOutX0, relOutY0, idx); |
|
scalar_t v = s_tileUpXY[idx] * (scalar_t)c_fd[0]; |
|
|
|
int outX = tileOutX + relOutX0; |
|
int outY = tileOutY + relOutY0; |
|
if ((uint32_t)outX < p.yShape.x && (uint32_t)outY < p.yShape.y) |
|
*((T*)((char*)p.y + (outX * get_stride<index_t>(p.yStride.x) + outY * get_stride<index_t>(p.yStride.y) + mapOfsOut))) = (T)v; |
|
} |
|
} |
|
} |
|
|
|
if (!enableXrep) |
|
break; |
|
} |
|
} |
|
|
|
|
|
|
|
|
|
|
|
template <class T, bool signWrite, bool signRead> |
|
static __global__ void filtered_lrelu_act_kernel(filtered_lrelu_act_kernel_params p) |
|
{ |
|
typedef typename InternalType<T>::scalar_t scalar_t; |
|
|
|
|
|
int32_t x = threadIdx.x + blockIdx.x * blockDim.x; |
|
int32_t ymax = signWrite ? p.sShape.y : p.xShape.y; |
|
int32_t qmax = p.xShape.z * p.xShape.w; |
|
|
|
|
|
for (int32_t q = blockIdx.z; q < qmax; q += gridDim.z) |
|
for (int32_t y = blockIdx.y; y < ymax; y += gridDim.y) |
|
{ |
|
|
|
int32_t w = q / p.xShape.z; |
|
int32_t z = q - w * p.xShape.z; |
|
|
|
|
|
if (signWrite) |
|
{ |
|
|
|
uint32_t s = 0; |
|
if (x < p.xShape.x && y < p.xShape.y) |
|
{ |
|
int64_t ix = x * p.xStride.x + y * p.xStride.y + z * p.xStride.z + w * p.xStride.w; |
|
T* pv = ((T*)p.x) + ix; |
|
scalar_t v = (scalar_t)(*pv); |
|
|
|
|
|
v *= p.gain; |
|
if (v < 0.f) |
|
{ |
|
v *= p.slope; |
|
s = 1; |
|
} |
|
if (fabsf(v) > p.clamp) |
|
{ |
|
v = InternalType<T>::clamp(v, p.clamp); |
|
s = 2; |
|
} |
|
|
|
*pv = (T)v; |
|
} |
|
|
|
|
|
uint32_t m = (threadIdx.x & 16) ? 0xffff0000u : 0x0000ffffu; |
|
s <<= ((threadIdx.x & 15) << 1); |
|
s |= __shfl_xor_sync(m, s, 1); |
|
s |= __shfl_xor_sync(m, s, 2); |
|
s |= __shfl_xor_sync(m, s, 4); |
|
s |= __shfl_xor_sync(m, s, 8); |
|
|
|
|
|
if (!(threadIdx.x & 15) && x < p.sShape.x) |
|
{ |
|
uint64_t is = x + p.sShape.x * (y + (int64_t)p.sShape.y * q); |
|
((uint32_t*)p.s)[is >> 4] = s; |
|
} |
|
} |
|
else if (signRead) |
|
{ |
|
|
|
if (x < p.xShape.x) |
|
{ |
|
int64_t ix = x * p.xStride.x + y * p.xStride.y + z * p.xStride.z + w * p.xStride.w; |
|
T* pv = ((T*)p.x) + ix; |
|
scalar_t v = (scalar_t)(*pv); |
|
v *= p.gain; |
|
|
|
|
|
uint32_t sx = x + p.sOfs.x; |
|
uint32_t sy = y + p.sOfs.y; |
|
|
|
|
|
if (sx < p.sShape.x && sy < p.sShape.y) |
|
{ |
|
uint64_t is = (sx >> 2) + (p.sShape.x >> 2) * (sy + (uint64_t)p.sShape.y * q); |
|
unsigned char s = p.s[is]; |
|
s >>= (sx & 3) << 1; |
|
if (s & 1) |
|
v *= p.slope; |
|
if (s & 2) |
|
v = 0.f; |
|
} |
|
|
|
*pv = (T)v; |
|
} |
|
} |
|
else |
|
{ |
|
|
|
if (x < p.xShape.x) |
|
{ |
|
int64_t ix = x * p.xStride.x + y * p.xStride.y + z * p.xStride.z + w * p.xStride.w; |
|
T* pv = ((T*)p.x) + ix; |
|
scalar_t v = (scalar_t)(*pv); |
|
v *= p.gain; |
|
if (v < 0.f) |
|
v *= p.slope; |
|
if (fabsf(v) > p.clamp) |
|
v = InternalType<T>::clamp(v, p.clamp); |
|
*pv = (T)v; |
|
} |
|
} |
|
} |
|
} |
|
|
|
template <class T, bool signWrite, bool signRead> void* choose_filtered_lrelu_act_kernel(void) |
|
{ |
|
return (void*)filtered_lrelu_act_kernel<T, signWrite, signRead>; |
|
} |
|
|
|
|
|
|
|
|
|
template <class T, class index_t, bool signWrite, bool signRead> filtered_lrelu_kernel_spec choose_filtered_lrelu_kernel(const filtered_lrelu_kernel_params& p, int sharedKB) |
|
{ |
|
filtered_lrelu_kernel_spec s = { 0 }; |
|
|
|
|
|
#define CASE(SH, U, FU, D, FD, MODE, TW, TH, W, XR, WS) \ |
|
if (sharedKB >= SH) \ |
|
if ((p.fuShape.y == 0 && (MODE == MODE_SUSD || MODE == MODE_SUFD)) || (p.fuShape.y > 0 && (MODE == MODE_FUSD || MODE == MODE_FUFD))) \ |
|
if ((p.fdShape.y == 0 && (MODE == MODE_SUSD || MODE == MODE_FUSD)) || (p.fdShape.y > 0 && (MODE == MODE_SUFD || MODE == MODE_FUFD))) \ |
|
if (p.up == U && p.fuShape.x <= FU && p.fuShape.y <= FU && p.down == D && p.fdShape.x <= FD && p.fdShape.y <= FD) \ |
|
{ \ |
|
static_assert((D*TW % 4) == 0, "down * tileWidth must be divisible by 4"); \ |
|
static_assert(FU % U == 0, "upscaling filter size must be multiple of upscaling factor"); \ |
|
static_assert(FD % D == 0, "downscaling filter size must be multiple of downscaling factor"); \ |
|
s.setup = (void*)setup_filters_kernel; \ |
|
s.exec = (void*)filtered_lrelu_kernel<T, index_t, SH, signWrite, signRead, MODE, U, FU, D, FD, TW, TH, W*32, !!XR, !!WS>; \ |
|
s.tileOut = make_int2(TW, TH); \ |
|
s.numWarps = W; \ |
|
s.xrep = XR; \ |
|
s.dynamicSharedKB = (SH == 48) ? 0 : SH; \ |
|
return s; \ |
|
} |
|
|
|
|
|
|
|
|
|
|
|
CASE(48, 1,1, 1,1, MODE_FUFD, 64, 178, 32, 0, 0) |
|
CASE(48, 2,8, 1,1, MODE_SUFD, 152, 95, 16, 0, 0) |
|
CASE(48, 1,1, 2,8, MODE_FUSD, 56, 22, 16, 0, 0) |
|
CASE(48, 2,8, 2,8, MODE_SUSD, 56, 29, 16, 11, 0) |
|
CASE(48, 2,8, 2,8, MODE_FUSD, 60, 28, 16, 0, 0) |
|
CASE(48, 2,8, 2,8, MODE_SUFD, 56, 28, 16, 0, 0) |
|
CASE(48, 4,16, 2,8, MODE_SUSD, 56, 31, 16, 11, 0) |
|
CASE(48, 4,16, 2,8, MODE_SUFD, 56, 36, 16, 0, 0) |
|
CASE(48, 2,8, 4,16, MODE_SUSD, 16, 22, 16, 12, 0) |
|
CASE(48, 2,8, 4,16, MODE_FUSD, 29, 15, 16, 0, 0) |
|
CASE(48, 2,12, 1,1, MODE_SUFD, 96, 150, 28, 0, 0) |
|
CASE(48, 1,1, 2,12, MODE_FUSD, 32, 35, 24, 0, 0) |
|
CASE(48, 2,12, 2,12, MODE_SUSD, 32, 46, 16, 10, 0) |
|
CASE(48, 2,12, 2,12, MODE_FUSD, 58, 28, 24, 8, 0) |
|
CASE(48, 2,12, 2,12, MODE_SUFD, 52, 28, 16, 0, 0) |
|
CASE(48, 4,24, 2,12, MODE_SUSD, 32, 51, 16, 5, 0) |
|
CASE(48, 4,24, 2,12, MODE_SUFD, 32, 56, 16, 6, 0) |
|
CASE(48, 2,12, 4,24, MODE_SUSD, 16, 18, 16, 12, 0) |
|
CASE(96, 2,12, 4,24, MODE_FUSD, 27, 31, 32, 6, 0) |
|
CASE(48, 2,12, 4,24, MODE_FUSD, 27, 13, 24, 0, 0) |
|
CASE(48, 2,16, 1,1, MODE_SUFD, 148, 89, 24, 0, 0) |
|
CASE(48, 1,1, 2,16, MODE_FUSD, 32, 31, 16, 5, 0) |
|
CASE(48, 2,16, 2,16, MODE_SUSD, 32, 41, 16, 9, 0) |
|
CASE(48, 2,16, 2,16, MODE_FUSD, 56, 26, 24, 0, 0) |
|
CASE(48, 2,16, 2,16, MODE_SUFD, 32, 40, 16, 0, 0) |
|
CASE(48, 4,32, 2,16, MODE_SUSD, 32, 46, 24, 5, 0) |
|
CASE(48, 4,32, 2,16, MODE_SUFD, 32, 50, 16, 0, 0) |
|
CASE(96, 2,16, 4,32, MODE_SUSD, 24, 24, 32, 12, 1) |
|
CASE(48, 2,16, 4,32, MODE_SUSD, 16, 13, 16, 10, 1) |
|
CASE(96, 2,16, 4,32, MODE_FUSD, 25, 28, 28, 4, 0) |
|
CASE(48, 2,16, 4,32, MODE_FUSD, 25, 10, 24, 0, 0) |
|
|
|
#undef CASE |
|
return s; |
|
} |
|
|
|
|
|
|