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
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
|
/* MATCHING eGPUTextureType. */
#define GPU_TEXTURE_1D (1 << 0)
#define GPU_TEXTURE_2D (1 << 1)
#define GPU_TEXTURE_3D (1 << 2)
#define GPU_TEXTURE_CUBE (1 << 3)
#define GPU_TEXTURE_ARRAY (1 << 4)
#define GPU_TEXTURE_BUFFER (1 << 5)
#define GPU_TEXTURE_1D_ARRAY (GPU_TEXTURE_1D | GPU_TEXTURE_ARRAY)
#define GPU_TEXTURE_2D_ARRAY (GPU_TEXTURE_2D | GPU_TEXTURE_ARRAY)
#define GPU_TEXTURE_CUBE_ARRAY (GPU_TEXTURE_CUBE | GPU_TEXTURE_ARRAY)
/* Determine input texture type. */
#if IS_DEPTH_FORMAT == 1
# define TEX_NAME_BASE depth
#else
# define TEX_NAME_BASE texture
#endif
#define JOIN(x, y) x##y
#define FUNC_NAME(x, y) JOIN(x, y)
/* Assign parameters based on texture type. */
#if TEX_TYPE == GPU_TEXTURE_1D
# define TEX_TYPE_NAME FUNC_NAME(TEX_NAME_BASE, 1d)
# define DIMS 1
#elif TEX_TYPE == GPU_TEXTURE_2D
# define TEX_TYPE_NAME FUNC_NAME(TEX_NAME_BASE, 2d)
# define DIMS 2
#elif TEX_TYPE == GPU_TEXTURE_3D
# define TEX_TYPE_NAME FUNC_NAME(TEX_NAME_BASE, 3d)
# define DIMS 3
#elif TEX_TYPE == GPU_TEXTURE_1D_ARRAY
# define TEX_TYPE_NAME FUNC_NAME(TEX_NAME_BASE, 1d_array)
# define DIMS 2
#elif TEX_TYPE == GPU_TEXTURE_2D_ARRAY
# define TEX_TYPE_NAME FUNC_NAME(TEX_NAME_BASE, 2d_array)
# define DIMS 3
#endif
/* Position dimensionality for threadgroup. */
#if DIMS == 1
# define POSITION_TYPE uint
#elif DIMS == 2
# define POSITION_TYPE uint2
#elif DIMS == 3
# define POSITION_TYPE uint3
#endif
using namespace metal;
template<typename T> T denormalize(float val)
{
return T(float(DEPTH_SCALE_FACTOR) * val);
};
template<> int denormalize<int>(float val)
{
return int((float(DEPTH_SCALE_FACTOR) * val - 1.0f) / 2.0f);
}
template<> uint denormalize<uint>(float val)
{
return uint(float(DEPTH_SCALE_FACTOR) * val);
}
template<typename T> T convert_type(float type)
{
return T(type);
}
template<> uchar convert_type<uchar>(float val)
{
return uchar(val * float(0xFF));
}
template<> uint convert_type<uint>(float val)
{
return uint(val * float(0xFFFFFFFFu));
}
struct TextureReadParams {
int mip_index;
int extent[3];
int offset[3];
};
#if IS_DEPTH_FORMAT == 1
constexpr sampler pixelSampler = sampler(coord::pixel, address::clamp_to_edge, filter::nearest);
#endif
kernel void compute_texture_read(constant TextureReadParams ¶ms [[buffer(0)]],
device OUTPUT_DATA_TYPE *output_data [[buffer(1)]],
#if IS_DEPTH_FORMAT == 1
TEX_TYPE_NAME<float, access::sample> read_tex [[texture(0)]],
#else
TEX_TYPE_NAME<INPUT_DATA_TYPE, access::read> read_tex
[[texture(0)]],
#endif
POSITION_TYPE position [[thread_position_in_grid]])
{
/* Read colour. */
vec<INPUT_DATA_TYPE, 4> read_colour;
/* 1D TEXTURE */
#if TEX_TYPE == GPU_TEXTURE_1D
/* xx, yy, layer determined by kernel invocation pattern */
uint xx = position[0];
int index = (xx)*COMPONENT_COUNT_OUTPUT;
read_colour = read_tex.read(uint(params.offset[0]) + uint(xx));
/* 2D TEXTURE */
#elif TEX_TYPE == GPU_TEXTURE_2D
/* xx, yy, layer determined by kernel invocation pattern */
uint xx = position[0];
uint yy = position[1];
int index = (yy * params.extent[0] + xx) * COMPONENT_COUNT_OUTPUT;
/* Read data */
# if IS_DEPTH_FORMAT == 1
output_data[index] = denormalize<OUTPUT_DATA_TYPE>(
read_tex.sample(pixelSampler, float2(params.offset[0], params.offset[1]) + float2(xx, yy)));
# else
read_colour = read_tex.read(uint2(params.offset[0], params.offset[1]) + uint2(xx, yy));
# endif
/* 3D TEXTURE */
#elif TEX_TYPE == GPU_TEXTURE_3D
/* xx, yy, layer determined by kernel invocation pattern */
uint xx = position[0];
uint yy = position[1];
uint zz = position[2];
int index = (zz * (params.extent[0] * params.extent[1]) + yy * params.extnt[0] + xx) *
COMPONENT_COUNT_INPUT;
read_colour = read_tex.read(uint3(params.offset[0], params.offset[1], params.offset[2]) +
uint3(xx, yy, zz));
/* 1D ARRAY TEXTURE */
#elif TEX_TYPE == GPU_TEXTURE_1D_ARRAY
/* xx, yy, layer determined by kernel invocation pattern */
uint xx = position[0];
uint layer = position[1];
int index = (layer * params.extent[0] + xx) * COMPONENT_COUNT_OUTPUT;
read_colour = read_tex.read(uint(params.offset[0]) + uint(xx), uint(params.offset[1]) + layer);
/* 2D ARRAY TEXTURE */
#elif TEX_TYPE == GPU_TEXTURE_2D_ARRAY
/* xx, yy, layer determined by kernel invocation pattern */
uint xx = position[0];
uint yy = position[1];
uint layer = position[2];
int index = (layer * (params.extent[0] * params.extent[1]) + yy * params.extent[0] + xx) *
COMPONENT_COUNT_INPUT;
/* Read data */
# if IS_DEPTH_FORMAT == 1
output_data[index] = denormalize<OUTPUT_DATA_TYPE>(
read_tex.sample(pixelSampler,
float2(params.offset[0], params.offset[1]) + float2(xx, yy),
uint(params.offset[2] + layer)));
# else
read_colour = read_tex.read(uint2(params.offset[0], params.offset[1]) + uint2(xx, yy),
uint(params.offset[2] + layer));
# endif
#endif
/* Output per-component colour data. */
#if IS_DEPTH_FORMAT != 1
/* Write data to block */
for (int i = 0; i < WRITE_COMPONENT_COUNT; i++) {
output_data[index + i] = convert_type<OUTPUT_DATA_TYPE>(read_colour[i]);
}
/* Fill in empty cells if more components are being read than exist */
for (int i = COMPONENT_COUNT_INPUT; i < COMPONENT_COUNT_OUTPUT; i++) {
output_data[index + i] = convert_type<OUTPUT_DATA_TYPE>(0);
}
#endif
}
|
