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
// Copyright (c) 2012-2017 VideoStitch SAS
// Copyright (c) 2018 stitchEm
__global__ void diskKernel(global_mem uint32_t* dst, unsigned width, unsigned height, float aX, float aY, float t,
uint32_t color) {
unsigned x = get_global_id_x();
unsigned y = get_global_id_y();
if (x < width && y < height) {
if ((aX - x) * (aX - x) + (aY - y) * (aY - y) < t) {
dst[width * y + x] = color;
}
}
}
__global__ void diskSourceKernel(surface_t dst, unsigned width, unsigned height, float aX, float aY, float t,
uint32_t color) {
unsigned x = get_global_id_x();
unsigned y = get_global_id_y();
if (x < width && y < height) {
if ((aX - x) * (aX - x) + (aY - y) * (aY - y) < t) {
surface_write_i(color, dst, x, y);
}
}
}
__global__ void lineKernel(global_mem uint32_t* dst, unsigned width, unsigned height, float aX, float aY, float bX,
float bY, float t, uint32_t color) {
unsigned x = get_global_id_x();
unsigned y = get_global_id_y();
if (x < width && y < height) {
float sqrLen = (aX - bX) * (aX - bX) + (aY - bY) * (aY - bY);
float p = (((float)x - aX) * (bX - aX) + ((float)y - aY) * (bY - aY)) / sqrLen;
float sqrDst;
if (p <= 0.0f) {
sqrDst = (aX - (float)x) * (aX - (float)x) + (aY - (float)y) * (aY - (float)y);
} else if (p >= 1.0f) {
sqrDst = (bX - (float)x) * (bX - (float)x) + (bY - (float)y) * (bY - (float)y);
} else {
float projX = aX + p * (bX - aX);
float projY = aY + p * (bY - aY);
sqrDst = (projX - (float)x) * (projX - (float)x) + (projY - (float)y) * (projY - (float)y);
}
if (sqrDst * 4.0f < t) {
dst[width * y + x] = color;
}
}
}
__global__ void lineSourceKernel(surface_t dst, unsigned width, unsigned height, float aX, float aY, float bX, float bY,
float t, uint32_t color) {
unsigned x = get_global_id_x();
unsigned y = get_global_id_y();
if (x < width && y < height) {
float sqrLen = (aX - bX) * (aX - bX) + (aY - bY) * (aY - bY);
float p = (((float)x - aX) * (bX - aX) + ((float)y - aY) * (bY - aY)) / sqrLen;
float sqrDst;
if (p <= 0.0f) {
sqrDst = (aX - (float)x) * (aX - (float)x) + (aY - (float)y) * (aY - (float)y);
} else if (p >= 1.0f) {
sqrDst = (bX - (float)x) * (bX - (float)x) + (bY - (float)y) * (bY - (float)y);
} else {
float projX = aX + p * (bX - aX);
float projY = aY + p * (bY - aY);
sqrDst = (projX - (float)x) * (projX - (float)x) + (projY - (float)y) * (projY - (float)y);
}
if (sqrDst * 4.0f < t) {
surface_write_i(color, dst, x, y);
}
}
}
#define CIRCLE_KERNEL_IMPL(name, moreTest, moreDesc) \
__global__ void name(global_mem uint32_t* dst, unsigned width, unsigned height, float centerX, float centerY, \
float innerSqrRadius, float outerSqrRadius, uint32_t color) { \
unsigned x = get_global_id_x(); \
unsigned y = get_global_id_y(); \
if (x < width && y < height) { \
float sqrDst = (centerX - (float)x) * (centerX - (float)x) + (centerY - (float)y) * (centerY - (float)y); \
if (innerSqrRadius < sqrDst && sqrDst < outerSqrRadius && (moreTest)) { \
dst[width * y + x] = color; \
} \
} \
} \
\
__global__ void name##Source(surface_t dst, unsigned width, unsigned height, float centerX, float centerY, \
float innerSqrRadius, float outerSqrRadius, uint32_t color) { \
unsigned x = get_global_id_x(); \
unsigned y = get_global_id_y(); \
if (x < width && y < height) { \
float sqrDst = (centerX - (float)x) * (centerX - (float)x) + (centerY - (float)y) * (centerY - (float)y); \
if (innerSqrRadius < sqrDst && sqrDst < outerSqrRadius && (moreTest)) { \
surface_write_i(color, dst, x, y); \
} \
} \
}
CIRCLE_KERNEL_IMPL(circleKernel, true, .)
CIRCLE_KERNEL_IMPL(circleTRKernel, x + 1 > centerX && y < centerY + 1, (top right quarter only))
CIRCLE_KERNEL_IMPL(circleBRKernel, x + 1 > centerX && y > centerY - 1, (bottom right quarter only))
CIRCLE_KERNEL_IMPL(circleTKernel, y < centerY, (top only))
CIRCLE_KERNEL_IMPL(circleBKernel, y > centerY, (bottom only))