1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
|
#include "pipeline.h"
#include <algorithm>
PyramidLevel make_level(int width, int height) {
PyramidLevel lvl{width, height};
hip_check(hipMalloc(&lvl.color_smax, sizeof(float4) * width * height));
hip_check(hipMalloc(&lvl.tensor, sizeof(float4) * width * height));
return lvl;
}
void free_level(PyramidLevel& lvl) {
hipFree(lvl.color_smax);
hipFree(lvl.tensor);
lvl.color_smax = nullptr;
lvl.tensor = nullptr;
}
Pyramid make_pyramid(int base_width, int base_height, int num_levels) {
Pyramid pyr;
pyr.levels.reserve(num_levels);
int w = base_width, h = base_height;
for (int i = 0; i < num_levels; ++i) {
pyr.levels.push_back(make_level(w, h));
w = std::max(1, w / 2);
h = std::max(1, h / 2);
}
return pyr;
}
void destroy_pyramid(Pyramid& pyr) {
for (auto& lvl : pyr.levels) free_level(lvl);
pyr.levels.clear();
}
/* Lanczos3 */
__device__ __forceinline__ float lanczos3(float x) {
const float a = 3.0f;
if (x == 0.0f) return 1.0f;
if (x <= -a || x >= a) return 0.0f;
const float px = M_PI * x;
return (a * sinf(px) * sinf(px / a)) / (px * px);
}
__device__ float4 lanczos3_sample_axis(
const float4* src, int src_len, int src_stride, int src_base,
int dst_idx, float scale)
{
float center = (dst_idx + 0.5f) * scale - 0.5f;
int icenter = static_cast<int>(floorf(center));
float fscale = fmaxf(scale, 1.0f);
int radius = static_cast<int>(ceilf(3.0f * fscale));
float4 sum = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
float wsum = 0.0f;
for (int i = icenter - radius; i <= icenter + radius; ++i) {
int c = min(max(i, 0), src_len - 1);
float w = lanczos3((i - center) / fscale);
float4 s = src[src_base + c * src_stride];
sum.x += s.x * w; sum.y += s.y * w;
sum.z += s.z * w; sum.w += s.w * w;
wsum += w;
}
if (wsum > 1e-6f) {
float inv = 1.0f / wsum;
sum.x *= inv; sum.y *= inv; sum.z *= inv; sum.w *= inv;
}
return sum;
}
__global__ void downsample_horizontal(
const float4* src, int src_w, int h,
float4* dst, int dst_w, float scale_x)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= dst_w || y >= h) return;
dst[y * dst_w + x] =
lanczos3_sample_axis(src, src_w, 1, y * src_w, x, scale_x);
}
__global__ void downsample_vertical(
const float4* src, int w, int src_h,
float4* dst, int dst_h, float scale_y)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= w || y >= dst_h) return;
dst[y * w + x] =
lanczos3_sample_axis(src, src_h, w, x, y, scale_y);
}
static void downsample_level(const PyramidLevel& src, PyramidLevel& dst) {
float scale_x = static_cast<float>(src.width) / dst.width;
float scale_y = static_cast<float>(src.height) / dst.height;
float4* tmp;
hip_check(hipMalloc(&tmp, sizeof(float4) * dst.width * src.height));
dim3 block(16, 16);
dim3 grid_h((dst.width + 15) / 16, (src.height + 15) / 16);
downsample_horizontal<<<grid_h, block>>>(
src.color_smax, src.width, src.height, tmp, dst.width, scale_x);
hip_check(hipGetLastError());
dim3 grid_v((dst.width + 15) / 16, (dst.height + 15) / 16);
downsample_vertical<<<grid_v, block>>>(
tmp, dst.width, src.height, dst.color_smax, dst.height, scale_y);
hip_check(hipGetLastError());
hipFree(tmp);
}
void build_pyramid(Pyramid& pyr) {
for (size_t i = 1; i < pyr.levels.size(); ++i)
downsample_level(pyr.levels[i - 1], pyr.levels[i]);
}
|