UltraGrid
UltraGrid Low-Latency High-Resolution Transmission System
Loading...
Searching...
No Matches
kernels.cu
Go to the documentation of this file.
1
8/*
9 * Copyright (c) 2024 CESNET
10 * All rights reserved.
11 *
12 * Redistribution and use in source and binary forms, with or without
13 * modification, is permitted provided that the following conditions
14 * are met:
15 *
16 * 1. Redistributions of source code must retain the above copyright
17 * notice, this list of conditions and the following disclaimer.
18 *
19 * 2. Redistributions in binary form must reproduce the above copyright
20 * notice, this list of conditions and the following disclaimer in the
21 * documentation and/or other materials provided with the distribution.
22 *
23 * 3. Neither the name of CESNET nor the names of its contributors may be
24 * used to endorse or promote products derived from this software without
25 * specific prior written permission.
26 *
27 * THIS SOFTWARE IS PROVIDED BY THE AUTHORS AND CONTRIBUTORS
28 * "AS IS" AND ANY EXPRESSED OR IMPLIED WARRANTIES, INCLUDING,
29 * BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY
30 * AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO
31 * EVENT SHALL THE AUTHORS OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
32 * INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
33 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
34 * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
35 * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
36 * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR
37 * OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE,
38 * EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
39 */
40
41#include "kernels.hpp"
42
43#include <cstdint>
44#include <cstdio>
45#include <cuda_runtime_api.h>
46
47extern volatile int log_level;
48#define LOG_LEVEL_DEBUG 7
49
50#define MEASURE_KERNEL_DURATION_START(stream) \
51 cudaEvent_t t0, t1; \
52 cudaEventCreate(&t0); \
53 cudaEventCreate(&t1); \
54 cudaEventRecord(t0, stream);
55#define MEASURE_KERNEL_DURATION_STOP(stream) \
56 cudaEventRecord(t1, stream); \
57 cudaEventSynchronize(t1); \
58 float elapsedTime = NAN; \
59 cudaEventElapsedTime(&elapsedTime, t0, t1); \
60 if (log_level >= LOG_LEVEL_DEBUG) { \
61 printf("%s elapsed time: %f ms\n", __func__, elapsedTime); \
62 } \
63 if (elapsedTime > 10.0) { \
64 fprintf( \
65 stderr, \
66 "Kernel in func %s duration %f > 10 ms, please report!\n", \
67 __func__, elapsedTime); \
68 }
69
70// ___ _____ ____ ___ __ ___ ___ ___ __
71// / _ \ / ___/ / / / ( _ ) ____\ \ / _ \ < / |_ | / /
72// / , _// (_ / /_ _// _ |/___/ > > / , _/ / / / __/ / /__
73// /_/|_| \___/ /_/ \___/ /_/ /_/|_| /_/ /____/ /____/
74
78template <typename load_t>
79__device__ static void
80rt48_to_r12l_compute_blk(const uint8_t *in, uint8_t *out)
81{
82 // load the data from in to src_u32
83 auto *in_t = (load_t *) in;
84 uint32_t src_u32[12];
85 for (unsigned i = 0; i < sizeof src_u32 / sizeof src_u32[0]; ++i) {
86 static_assert(sizeof(load_t) == 2 || sizeof(load_t) == 4,
87 "Just uint{16,32} supported!");
88 if (sizeof(load_t) == 4) {
89 src_u32[i] = in_t[i];
90 } else {
91 src_u32[i] = in_t[2 * i] | in_t[2 * i + 1] << 16;
92 }
93 }
94
95 uint32_t dst_u32[9];
96 auto *dst = (uint8_t *) dst_u32;
97 auto *src = (uint8_t *) src_u32;
98
99 // 0
100 dst[0] = src[0] >> 4;
101 dst[0] |= src[1] << 4;
102 dst[1] = src[1] >> 4;
103 src += 2;
104
105 dst[1] |= src[0] & 0xF0;
106 dst[2] = src[1];
107 src += 2;
108
109 dst[3] = src[0] >> 4;
110 dst[3] |= src[1] << 4;
111 dst[4 + 0] = src[1] >> 4;
112 src += 2;
113
114 // 1
115 dst[4 + 0] |= src[0] & 0xF0;
116 dst[4 + 1] = src[1];
117 src += 2;
118
119 dst[4 + 2] = src[0] >> 4;
120 dst[4 + 2] |= src[1] << 4;
121 dst[4 + 3] = src[1] >> 4;
122 src += 2;
123
124 dst[4 + 3] |= src[0] & 0xF0;
125 dst[8 + 0] = src[1];
126 src += 2;
127
128 // 2
129 dst[8 + 1] = src[0] >> 4;
130 dst[8 + 1] |= src[1] << 4;
131 dst[8 + 2] = src[1] >> 4;
132 src += 2;
133
134 dst[8 + 2] |= src[0] & 0xF0;
135 dst[8 + 3] = src[1];
136 src += 2;
137
138 dst[12 + 0] = src[0] >> 4;
139 dst[12 + 0] |= src[1] << 4;
140 dst[12 + 1] = src[1] >> 4;
141 src += 2;
142
143 // 3
144 dst[12 + 1] |= src[0] & 0xF0;
145 dst[12 + 2] = src[1];
146 src += 2;
147
148 dst[12 + 3] = src[0] >> 4;
149 dst[12 + 3] |= src[1] << 4;
150 dst[16 + 0] = src[1] >> 4;
151 src += 2;
152
153 dst[16 + 0] |= src[0] & 0xF0;
154 dst[16 + 1] = src[1];
155 src += 2;
156
157 // 4
158 dst[16 + 2] = src[0] >> 4;
159 dst[16 + 2] |= src[1] << 4;
160 dst[16 + 3] = src[1] >> 4;
161 src += 2;
162
163 dst[16 + 3] |= src[0] & 0xF0;
164 dst[20 + 0] = src[1];
165 src += 2;
166
167 dst[20 + 1] = src[0] >> 4;
168 dst[20 + 1] |= src[1] << 4;
169 dst[20 + 2] = src[1] >> 4;
170 src += 2;
171
172 // 5
173 dst[20 + 2] |= src[0] & 0xF0;
174 dst[20 + 3] = src[1];
175 src += 2;
176
177 dst[24 + 0] = src[0] >> 4;
178 dst[24 + 0] |= src[1] << 4;
179 dst[24 + 1] = src[1] >> 4;
180 src += 2;
181
182 dst[24 + 1] |= src[0] & 0xF0;
183 dst[24 + 2] = src[1];
184 src += 2;
185
186 // 6
187 dst[24 + 3] = src[0] >> 4;
188 dst[24 + 3] |= src[1] << 4;
189 dst[28 + 0] = src[1] >> 4;
190 src += 2;
191
192 dst[28 + 0] |= src[0] & 0xF0;
193 dst[28 + 1] = src[1];
194 src += 2;
195
196 dst[28 + 2] = src[0] >> 4;
197 dst[28 + 2] |= src[1] << 4;
198 dst[28 + 3] = src[1] >> 4;
199 src += 2;
200
201 // 7
202 dst[28 + 3] |= src[0] & 0xF0;
203 dst[32 + 0] = src[1];
204 src += 2;
205
206 dst[32 + 1] = src[0] >> 4;
207 dst[32 + 1] |= src[1] << 4;
208 dst[32 + 2] = src[1] >> 4;
209 src += 2;
210
211 dst[32 + 2] |= src[0] & 0xF0;
212 dst[32 + 3] = src[1];
213 src += 2;
214
215 // store the result
216 auto *out_u32 = (uint32_t *) out;
217 for (unsigned i = 0; i < sizeof dst_u32 / sizeof dst_u32[0]; ++i) {
218 out_u32[i] = dst_u32[i];
219 }
220}
221
222template <typename load_t>
223__device__ static void
224rt48_to_r12l_compute_last_blk(uint8_t *src, uint8_t *dst, unsigned width)
225{
226 alignas(uint32_t) uint8_t tmp[48];
227 for (unsigned i = 0; i < width * 6; ++i) {
228 tmp[i] = src[i];
229 }
230 rt48_to_r12l_compute_blk<load_t>(tmp, dst);
231}
232
236template <typename load_t>
237__global__ static void
238kernel_rg48_to_r12l(uint8_t *in, uint8_t *out, unsigned size_x)
239{
240 unsigned position_x = threadIdx.x + blockIdx.x * blockDim.x;
241 unsigned position_y = threadIdx.y + blockIdx.y * blockDim.y;
242 if (position_x >= (size_x + 7) / 8) {
243 return;
244 }
245 uint8_t *src = in + 2 * (position_y * 3 * size_x + position_x * 3 * 8);
246 uint8_t *dst =
247 out + (position_y * ((size_x + 7) / 8) + position_x) * 36;
248
249 // handle incomplete blocks
250 if (position_x == size_x / 8) {
251 rt48_to_r12l_compute_last_blk<load_t>(src, dst,
252 size_x - position_x * 8);
253 return;
254 }
255 rt48_to_r12l_compute_blk<load_t>(src, dst);
256}
257
261int postprocess_rg48_to_r12l(
262 void * /* postprocessor */,
263 void * /* img_custom_data*/,
264 size_t /* img_custom_data_size */,
265 int size_x,
266 int size_y,
267 struct cmpto_j2k_dec_comp_format * /* comp_formats */,
268 int /* comp_count */,
269 void *input_samples,
270 size_t /* input_samples_size */,
271 void * /* temp_buffer */,
272 size_t /* temp_buffer_size */,
273 void * output_buffer,
274 size_t /* output_buffer_size */,
275 void * vstream
276) {
277 cudaStream_t stream = (cudaStream_t) vstream;
278 dim3 threads_per_block(256);
279 dim3 blocks((((size_x + 7) / 8) + 255) / 256, size_y);
280
281 MEASURE_KERNEL_DURATION_START(stream)
282
283 if (size_x % 2 == 0) {
284 kernel_rg48_to_r12l<uint32_t>
285 <<<blocks, threads_per_block, 0, (cudaStream_t) stream>>>(
286 (uint8_t *) input_samples, (uint8_t *) output_buffer,
287 size_x);
288 } else {
289 thread_local bool warn_print;
290 if (!warn_print) {
291 fprintf(stderr,
292 "%s: Odd width %d px will use slower kernel!\n",
293 __func__, size_x);
294 warn_print = true;
295 }
296 kernel_rg48_to_r12l<uint16_t>
297 <<<blocks, threads_per_block, 0, (cudaStream_t) stream>>>(
298 (uint8_t *) input_samples, (uint8_t *) output_buffer,
299 size_x);
300 }
301
302 MEASURE_KERNEL_DURATION_STOP(stream)
303
304 return cudaGetLastError();
305}
306
307// ___ ___ ___ __ __ ___ _____ ____ ___
308// / _ \ < / |_ | / / ____\ \ / _ \ / ___/ / / / ( _ )
309// / , _/ / / / __/ / /__/___/ > > / , _// (_ / /_ _// _ |
310// /_/|_| /_/ /____/ /____/ /_/ /_/|_| \___/ /_/ \___/
311
312template <typename store_t>
313__device__ static void r12l_to_rg48_compute_blk(const uint8_t *src,
314 uint8_t *dst);
315
316template <typename store_t>
317__global__ static void
318kernel_r12l_to_rg48(uint8_t *in, uint8_t *out, unsigned size_x)
319{
320 unsigned position_x = threadIdx.x + blockIdx.x * blockDim.x;
321 unsigned position_y = threadIdx.y + blockIdx.y * blockDim.y;
322 if (position_x >= (size_x + 7) / 8) {
323 return;
324 }
325 uint8_t *dst = out + 2 * (position_y * 3 * size_x + position_x * 3 * 8);
326 uint8_t *src =
327 in + (position_y * ((size_x + 7) / 8) + position_x) * 36;
328
329 if (position_x == size_x / 8) {
330 // compute the last incomplete block
331 alignas(uint32_t) uint8_t tmp[48];
332 r12l_to_rg48_compute_blk<store_t>(src, tmp);
333 for (unsigned i = 0; i < (size_x - position_x * 8) * 6; ++i) {
334 dst[i] = tmp[i];
335 }
336 return;
337 }
338 r12l_to_rg48_compute_blk<store_t>(src, dst);
339}
340
342template <typename store_t>
343__device__ static void
344r12l_to_rg48_compute_blk(const uint8_t *in, uint8_t *out)
345{
346 // load the data from in to src_u32
347 auto *in_u32 = (uint32_t *) in;
348 uint32_t src_u32[9];
349 for (unsigned i = 0; i < sizeof src_u32 / sizeof src_u32[0]; ++i) {
350 src_u32[i] = in_u32[i];
351 }
352
353 uint32_t dst_u32[12];
354 uint8_t *dst = (uint8_t *) dst_u32;
355 uint8_t *src = (uint8_t *) src_u32;
356
357 // 0
358 // R
359 *dst++ = src[0] << 4;
360 *dst++ = (src[1] << 4) | (src[0] >> 4);
361 // G
362 *dst++ = src[1] & 0xF0;
363 *dst++ = src[2];
364 // B
365 *dst++ = src[3] << 4;
366 *dst++ = (src[4 + 0] << 4) | (src[3] >> 4);
367
368 // 1
369 *dst++ = src[4 + 0] & 0xF0;
370 *dst++ = src[4 + 1];
371
372 *dst++ = src[4 + 2] << 4;
373 *dst++ = (src[4 + 3] << 4) | (src[4 + 2] >> 4);
374
375 *dst++ = src[4 + 3] & 0xF0;
376 *dst++ = src[8 + 0];
377
378 // 2
379 *dst++ = src[8 + 1] << 4;
380 *dst++ = (src[8 + 2] << 4) | (src[8 + 1] >> 4);
381
382 *dst++ = src[8 + 2] & 0xF0;
383 *dst++ = src[8 + 3];
384
385 *dst++ = src[12 + 0] << 4;
386 *dst++ = (src[12 + 1] << 4) | (src[12 + 0] >> 4);
387
388 // 3
389 *dst++ = src[12 + 1] & 0xF0;
390 *dst++ = src[12 + 2];
391
392 *dst++ = src[12 + 3] << 4;
393 *dst++ = (src[16 + 0] << 4) | (src[12 + 3] >> 4);
394
395 *dst++ = src[16 + 0] & 0xF0;
396 *dst++ = src[16 + 1];
397
398 // 4
399 *dst++ = src[16 + 2] << 4;
400 *dst++ = (src[16 + 3] << 4) | (src[16 + 2] >> 4);
401
402 *dst++ = src[16 + 3] & 0xF0;
403 *dst++ = src[20 + 0];
404
405 *dst++ = src[20 + 1] << 4;
406 *dst++ = (src[20 + 2] << 4) | (src[20 + 1] >> 4);
407
408 // 5
409 *dst++ = src[20 + 2] & 0xF0;
410 *dst++ = src[20 + 3];
411
412 *dst++ = src[24 + 0] << 4;
413 *dst++ = (src[24 + 1] << 4) | (src[24 + 0] >> 4);
414
415 *dst++ = src[24 + 1] & 0xF0;
416 *dst++ = src[24 + 2];
417
418 // 6
419 *dst++ = src[24 + 3] << 4;
420 *dst++ = (src[28 + 0] << 4) | (src[24 + 3] >> 4);
421
422 *dst++ = src[28 + 0] & 0xF0;
423 *dst++ = src[28 + 1];
424
425 *dst++ = src[28 + 2] << 4;
426 *dst++ = (src[28 + 3] << 4) | (src[28 + 2] >> 4);
427
428 // 7
429 *dst++ = src[28 + 3] & 0xF0;
430 *dst++ = src[32 + 0];
431
432 *dst++ = src[32 + 1] << 4;
433 *dst++ = (src[32 + 2] << 4) | (src[32 + 1] >> 4);
434
435 *dst++ = src[32 + 2] & 0xF0;
436 *dst++ = src[32 + 3];
437
438 // store the result
439 auto *out_t = (store_t *) out;
440 for (unsigned i = 0; i < sizeof dst_u32 / sizeof dst_u32[0]; ++i) {
441 static_assert(sizeof(store_t) == 2 || sizeof(store_t) == 4,
442 "Just uint{16,32} supported!");
443 if (sizeof(store_t) == 4) {
444 out_t[i] = dst_u32[i];
445 } else {
446 out_t[2 * i] = dst_u32[i] & 0xFFFFU;
447 out_t[2 * i + 1] = dst_u32[i] >> 16;
448 }
449 }
450}
451
452int
453preprocess_r12l_to_rg48(void *preprocessor, void *img_custom_data,
454 size_t img_custom_data_size, int size_x, int size_y,
455 struct cmpto_j2k_enc_comp_format *comp_formats,
456 int comp_count, void *input_samples,
457 size_t input_samples_size, void *output_samples,
458 size_t output_samples_size, void *vstream)
459{
460 (void) preprocessor, (void) img_custom_data,
461 (void) img_custom_data_size, (void) comp_formats, (void) comp_count,
462 (void) input_samples_size, (void) output_samples_size;
463
464 cudaStream_t stream = (cudaStream_t) vstream;
465 dim3 threads_per_block(256);
466 dim3 blocks((((size_x+ 7) / 8) + 255) / 256, size_y);
467
468 MEASURE_KERNEL_DURATION_START(stream)
469 if (size_x % 2 == 0) {
470 kernel_r12l_to_rg48<uint32_t>
471 <<<blocks, threads_per_block, 0, stream>>>(
472 (uint8_t *) input_samples, (uint8_t *) output_samples,
473 size_x);
474 } else {
475 thread_local bool warn_print;
476 if (!warn_print) {
477 fprintf(stderr,
478 "%s: Odd width %d px will use slower kernel!\n",
479 __func__, size_x);
480 warn_print = true;
481 }
482 kernel_r12l_to_rg48<uint16_t>
483 <<<blocks, threads_per_block, 0, stream>>>(
484 (uint8_t *) input_samples, (uint8_t *) output_samples,
485 size_x);
486 }
487 MEASURE_KERNEL_DURATION_STOP(stream)
488
489 return cudaGetLastError();
490}
cudaGetLastError
#define cudaGetLastError
Definition cuda_runtime.h:46
kernel_rg48_to_r12l
static __global__ void kernel_rg48_to_r12l(uint8_t *in, uint8_t *out, unsigned size_x)
Definition kernels.cu:238
log_level
volatile int log_level
Definition debug.cpp:65
rt48_to_r12l_compute_blk
static __device__ void rt48_to_r12l_compute_blk(const uint8_t *in, uint8_t *out)
Definition kernels.cu:80
r12l_to_rg48_compute_blk
static __device__ void r12l_to_rg48_compute_blk(const uint8_t *src, uint8_t *dst)
adapted variant of vc_copylineR12LtoRG48
Definition kernels.cu:344
preprocess_r12l_to_rg48
int preprocess_r12l_to_rg48(void *preprocessor, void *img_custom_data, size_t img_custom_data_size, int size_x, int size_y, struct cmpto_j2k_enc_comp_format *comp_formats, int comp_count, void *input_samples, size_t input_samples_size, void *output_samples, size_t output_samples_size, void *vstream)
Definition kernels.cu:453
MEASURE_KERNEL_DURATION_START
#define MEASURE_KERNEL_DURATION_START(stream)
Definition kernels.cu:50
postprocess_rg48_to_r12l
int postprocess_rg48_to_r12l(void *, void *, size_t, int size_x, int size_y, struct cmpto_j2k_dec_comp_format *, int, void *input_samples, size_t, void *, size_t, void *output_buffer, size_t, void *vstream)
Definition kernels.cu:261
MEASURE_KERNEL_DURATION_STOP
#define MEASURE_KERNEL_DURATION_STOP(stream)
Definition kernels.cu:55
kernel_r12l_to_rg48
static __global__ void kernel_r12l_to_rg48(uint8_t *in, uint8_t *out, unsigned size_x)
Definition kernels.cu:318
rt48_to_r12l_compute_last_blk
static __device__ void rt48_to_r12l_compute_last_blk(uint8_t *src, uint8_t *dst, unsigned width)
Definition kernels.cu:224
Generated on Wed Sep 25 2024 14:27:16 for UltraGrid by doxygen 1.12.0