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
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
// Copyright (c) 2012-2017 VideoStitch SAS
// Copyright (c) 2018 stitchEm
#include "../gpuKernelDef.h"
#include "../image/imageFormat.h"
static inline bool OutputRectCropper_isPanoPointVisible(int x, int y, int panoWidth, int panoHeight) { return true; }
static inline bool OutputCircleCropper_isPanoPointVisible(int x, int y, int panoWidth, int panoHeight) { return true; }
static inline bool isWithinCropRect(const float2 uv, float width, float height, float cLeft, float cRight, float cTop,
float cBottom) {
return 0.0f <= uv.x && uv.x < width && 0.0f <= uv.y && uv.y < height && cLeft <= uv.x && uv.x <= cRight &&
cTop <= uv.y && uv.y <= cBottom;
}
static inline bool isWithinCropCircle(const float2 uv, float width, float height, float cLeft, float cRight, float cTop,
float cBottom) {
const float centerX = (cRight + cLeft) / 2.0f;
const float centerY = (cBottom + cTop) / 2.0f;
const float radius = fmin(cRight - cLeft, cBottom - cTop) / 2.0f;
return 0.0f <= uv.x && uv.x < width && 0.0f <= uv.y && uv.y < height &&
(uv.x - centerX) * (uv.x - centerX) + (uv.y - centerY) * (uv.y - centerY) < radius * radius;
}
#include "backend/common/core/types.hpp"
#include "mapFunction.h"
#include "backend/common/core1/zoneKernel.gpu"
static const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
#define OUTPUTPROJECTION ErectToSphere
#define OUTPUTCROPPER OutputRectCropper_isPanoPointVisible
#define REMAPKERNEL remap_equirectangular
#include "remapKernel.cl.incl"
#undef REMAPKERNEL
#undef OUTPUTCROPPER
#undef OUTPUTPROJECTION
#define OUTPUTPROJECTION RectToSphere
#define OUTPUTCROPPER OutputRectCropper_isPanoPointVisible
#define REMAPKERNEL remap_rectilinear
#include "remapKernel.cl.incl"
#undef REMAPKERNEL
#undef OUTPUTCROPPER
#undef OUTPUTPROJECTION
#define OUTPUTPROJECTION FisheyeToSphere
#define OUTPUTCROPPER OutputRectCropper_isPanoPointVisible
#define REMAPKERNEL remap_fullframe_fisheye
#include "remapKernel.cl.incl"
#undef REMAPKERNEL
#undef OUTPUTCROPPER
#undef OUTPUTPROJECTION
#define OUTPUTPROJECTION FisheyeToSphere
#define OUTPUTCROPPER OutputCircleCropper_isPanoPointVisible
#define REMAPKERNEL remap_circular_fisheye
#include "remapKernel.cl.incl"
#undef REMAPKERNEL
#undef OUTPUTCROPPER
#undef OUTPUTPROJECTION
#define OUTPUTPROJECTION StereoToSphere
#define OUTPUTCROPPER OutputRectCropper_isPanoPointVisible
#define REMAPKERNEL remap_stereographic
#include "remapKernel.cl.incl"
#undef REMAPKERNEL
#undef OUTPUTCROPPER
#undef OUTPUTPROJECTION
kernel void remap_cubemap(global uint32_t* g_odata, read_only image2d_t tex, int panoWidth, int panoHeight,
const float2 panoScale, const vsfloat3x3 R, int faceDim, global uint32_t* xPositive,
global uint32_t* xNegative, global uint32_t* yPositive, global uint32_t* yNegative,
global uint32_t* zPositive, global uint32_t* zNegative) {
const int x = (int)get_global_id(0);
const int y = (int)get_global_id(1);
if (x < faceDim && y < faceDim) {
/* compensate fetching offset with CLK_FILTER_LINEAR by adding 0.5f */
float2 uv = make_float2(x + 0.5f, y + 0.5f);
uv = (uv / faceDim) * 2.f - make_float2(1.f, 1.f);
float3 pt = {0, 0, 0};
for (unsigned int face = 0; face < 6; face++) {
// Layer 0 is positive X face
if (face == 0) {
pt.x = 1;
pt.y = -uv.y;
pt.z = -uv.x;
}
// Layer 1 is negative X face
else if (face == 1) {
pt.x = -1;
pt.y = -uv.y;
pt.z = uv.x;
}
// Layer 2 is positive Y face
else if (face == 2) {
pt.x = uv.x;
pt.y = 1;
pt.z = uv.y;
}
// Layer 3 is negative Y face
else if (face == 3) {
pt.x = uv.x;
pt.y = -1;
pt.z = -uv.y;
}
// Layer 4 is positive Z face
else if (face == 4) {
pt.x = uv.x;
pt.y = -uv.y;
pt.z = 1;
}
// Layer 5 is negative Z face
else if (face == 5) {
pt.x = -uv.x;
pt.y = -uv.y;
pt.z = -1;
}
pt = rotateSphere(pt, R);
float2 xy = SphereToErect(pt);
xy *= panoScale;
/**
* See notes in warp kernel
*/
xy.x += panoWidth / 2.0f;
xy.y += panoHeight / 2.0f;
const float4 im = read_imagef(tex, smp, xy) * 255.0f;
if (face == 0) {
xPositive[y * faceDim + x] = (uchar)im.s3 << 24 | (uchar)im.s2 << 16 | (uchar)im.s1 << 8 | (uchar)im.s0;
} else if (face == 1) {
xNegative[y * faceDim + x] = (uchar)im.s3 << 24 | (uchar)im.s2 << 16 | (uchar)im.s1 << 8 | (uchar)im.s0;
} else if (face == 2) {
yPositive[y * faceDim + x] = (uchar)im.s3 << 24 | (uchar)im.s2 << 16 | (uchar)im.s1 << 8 | (uchar)im.s0;
} else if (face == 3) {
yNegative[y * faceDim + x] = (uchar)im.s3 << 24 | (uchar)im.s2 << 16 | (uchar)im.s1 << 8 | (uchar)im.s0;
} else if (face == 4) {
zPositive[y * faceDim + x] = (uchar)im.s3 << 24 | (uchar)im.s2 << 16 | (uchar)im.s1 << 8 | (uchar)im.s0;
} else if (face == 5) {
zNegative[y * faceDim + x] = (uchar)im.s3 << 24 | (uchar)im.s2 << 16 | (uchar)im.s1 << 8 | (uchar)im.s0;
}
}
}
}