kernels.cu File Reference

#include "kernels.hpp"
#include <cstdint>
#include <cstdio>
#include <cuda_runtime_api.h>

Include dependency graph for kernels.cu:

Go to the source code of this file.

Macros
#define	LOG_LEVEL_DEBUG   7
#define	MEASURE_KERNEL_DURATION_START(stream)
#define	MEASURE_KERNEL_DURATION_STOP(stream)

Functions
template<typename store_t >
static __global__ void	kernel_r12l_to_rg48 (uint8_t *in, uint8_t *out, unsigned size_x)
template<typename load_t >
static __global__ void	kernel_rg48_to_r12l (uint8_t *in, uint8_t *out, unsigned size_x)
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)
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)
template<typename store_t >
static __device__ void	r12l_to_rg48_compute_blk (const uint8_t *in, uint8_t *out)
	adapted variant of vc_copylineR12LtoRG48
template<typename load_t >
static __device__ void	rt48_to_r12l_compute_blk (const uint8_t *in, uint8_t *out)
template<typename load_t >
static __device__ void	rt48_to_r12l_compute_last_blk (uint8_t *src, uint8_t *dst, unsigned width)

Variables
volatile int	log_level

Detailed Description

Author

Martin Pulec pulec.nosp@m.@ces.nosp@m.net.c.nosp@m.z

This file hosts various CUDA kernels. Currently there are only kernels for cmpto_j2k compression and decompression.

Definition in file kernels.cu.

Macro Definition Documentation

LOG_LEVEL_DEBUG

#define LOG_LEVEL_DEBUG   7

Definition at line 48 of file kernels.cu.

Referenced by _dprintf(), audio_buffer_read(), audio_update_recv_buf(), blacklist_current_out_codec(), SampleGrabberCallback::BufferCB(), vidcap_state_aja::CaptureFrames(), configure_with(), debug_dump(), decompress_thread(), display_decklink_put_audio_frame(), display_sdl_run(), image_pattern_bars::fill(), image_pattern_ebu_smpte_bars< f, columns >::fill(), image_pattern_smpte_bars::fill(), gl_process_frames(), glsl_compile_link(), init_rtsp(), juice_log_handler(), libavcodec_decompress(), main(), on_process(), parallel_conv(), parse_log_cfg(), print_packet_info(), print_rtsp_header(), probe_internal(), probe_internal_cineform(), ultragrid::aja::display::process_frames(), process_video_pkt(), rg48_to_r12l(), rtp_init_if(), rtp_init_with_udp_socket(), rtsp_get_parameters(), ScreenCastPortal_impl::session_created(), set_display_mode_properties(), setup_nat_pmp(), signal_handler(), ScreenCastPortal_impl::started(), tfrc_done(), AudioDriftFixer::update(), vidcap_dshow_grab(), vidcap_file_worker(), vidcap_pw_init(), vidcap_screen_pw_done(), and videotoolbox_init().

MEASURE_KERNEL_DURATION_START

#define MEASURE_KERNEL_DURATION_START

(

stream

)

Value:

        cudaEvent_t t0, t1; \
        cudaEventCreate(&t0); \
        cudaEventCreate(&t1); \
        cudaEventRecord(t0, stream);

Definition at line 50 of file kernels.cu.

#define MEASURE_KERNEL_DURATION_START(stream) \
        cudaEvent_t t0, t1; \
        cudaEventCreate(&t0); \
        cudaEventCreate(&t1); \
        cudaEventRecord(t0, stream);

Referenced by postprocess_rg48_to_r12l(), and preprocess_r12l_to_rg48().

MEASURE_KERNEL_DURATION_STOP

#define MEASURE_KERNEL_DURATION_STOP

(

stream

)

Value:

        cudaEventRecord(t1, stream); \
        cudaEventSynchronize(t1); \
        float elapsedTime = NAN; \
        cudaEventElapsedTime(&elapsedTime, t0, t1); \
        if (log_level >= LOG_LEVEL_DEBUG) { \
                printf("%s elapsed time: %f ms\n", __func__, elapsedTime); \
        } \
        if (elapsedTime > 10.0) { \
                fprintf( \
                    stderr, \
                    "Kernel in func %s duration %f > 10 ms, please report!\n", \
                    __func__, elapsedTime); \
        }

Definition at line 55 of file kernels.cu.

#define MEASURE_KERNEL_DURATION_STOP(stream) \
        cudaEventRecord(t1, stream); \
        cudaEventSynchronize(t1); \
        float elapsedTime = NAN; \
        cudaEventElapsedTime(&elapsedTime, t0, t1); \
        if (log_level >= LOG_LEVEL_DEBUG) { \
                printf("%s elapsed time: %f ms\n", __func__, elapsedTime); \
        } \
        if (elapsedTime > 10.0) { \
                fprintf( \
                    stderr, \
                    "Kernel in func %s duration %f > 10 ms, please report!\n", \
                    __func__, elapsedTime); \
        }

Referenced by postprocess_rg48_to_r12l(), and preprocess_r12l_to_rg48().

Function Documentation

kernel_r12l_to_rg48()

template<typename store_t >

static __global__ void kernel_r12l_to_rg48 ( uint8_t * in, uint8_t * out, unsigned size_x )

static

Definition at line 318 of file kernels.cu.

{
        unsigned position_x = threadIdx.x + blockIdx.x * blockDim.x;
        unsigned position_y = threadIdx.y + blockIdx.y * blockDim.y;
        if (position_x >= (size_x + 7) / 8) {
                return;
        }
        uint8_t *dst = out + 2 * (position_y * 3 * size_x + position_x * 3 * 8);
        uint8_t *src =
            in + (position_y * ((size_x + 7) / 8) + position_x) * 36;
 
        if (position_x == size_x / 8) {
                // compute the last incomplete block
                alignas(uint32_t) uint8_t tmp[48];
                r12l_to_rg48_compute_blk<store_t>(src, tmp);
                for (unsigned i = 0; i < (size_x - position_x * 8) * 6; ++i) {
                        dst[i] = tmp[i];
                }
                return;
        }
        r12l_to_rg48_compute_blk<store_t>(src, dst);
}

References r12l_to_rg48_compute_blk().

Referenced by preprocess_r12l_to_rg48().

Here is the call graph for this function:

kernel_rg48_to_r12l()

template<typename load_t >

static __global__ void kernel_rg48_to_r12l ( uint8_t * in, uint8_t * out, unsigned size_x )

static

Todo

fix the last block for widths not divisible by 8

Definition at line 238 of file kernels.cu.

{
        unsigned position_x = threadIdx.x + blockIdx.x * blockDim.x;
        unsigned position_y = threadIdx.y + blockIdx.y * blockDim.y;
        if (position_x >= (size_x + 7) / 8) {
                return;
        }
        uint8_t *src = in + 2 * (position_y * 3 * size_x + position_x * 3 * 8);
        uint8_t *dst =
            out + (position_y * ((size_x + 7) / 8) + position_x) * 36;
 
        // handle incomplete blocks
        if (position_x == size_x / 8) {
                rt48_to_r12l_compute_last_blk<load_t>(src, dst,
                                                      size_x - position_x * 8);
                return;
        }
        rt48_to_r12l_compute_blk<load_t>(src, dst);
}

References rt48_to_r12l_compute_blk(), and rt48_to_r12l_compute_last_blk().

Referenced by postprocess_rg48_to_r12l().

Here is the call graph for this function:

postprocess_rg48_to_r12l()

int postprocess_rg48_to_r12l	(	void *	postprocessor,
		void *	img_custom_data,
		size_t	img_custom_data_size,
		int	size_x,
		int	size_y,
		struct cmpto_j2k_dec_comp_format *	comp_formats,
		int	comp_count,
		void *	input_samples,
		size_t	input_samples_size,
		void *	temp_buffer,
		size_t	temp_buffer_size,
		void *	output_buffer,
		size_t	output_buffer_size,
		void *	vstream )

See also

cmpto_j2k_dec_postprocessor_run_callback_cuda

Definition at line 261 of file kernels.cu.

  {
        cudaStream_t stream = (cudaStream_t) vstream;
        dim3 threads_per_block(256);
        dim3 blocks((((size_x + 7) / 8) + 255) / 256, size_y);
 
        MEASURE_KERNEL_DURATION_START(stream)
 
        if (size_x % 2 == 0) {
                kernel_rg48_to_r12l<uint32_t>
                    <<<blocks, threads_per_block, 0, (cudaStream_t) stream>>>(
                        (uint8_t *) input_samples, (uint8_t *) output_buffer,
                        size_x);
        } else {
                thread_local bool warn_print;
                if (!warn_print) {
                        fprintf(stderr,
                                "%s: Odd width %d px will use slower kernel!\n",
                                __func__, size_x);
                        warn_print = true;
                }
                kernel_rg48_to_r12l<uint16_t>
                    <<<blocks, threads_per_block, 0, (cudaStream_t) stream>>>(
                        (uint8_t *) input_samples, (uint8_t *) output_buffer,
                        size_x);
        }
 
        MEASURE_KERNEL_DURATION_STOP(stream)
 
        return cudaGetLastError();
}

References cudaGetLastError, kernel_rg48_to_r12l(), MEASURE_KERNEL_DURATION_START, and MEASURE_KERNEL_DURATION_STOP.

Here is the call graph for this function:

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 at line 453 of file kernels.cu.

{
        (void) preprocessor, (void) img_custom_data,
            (void) img_custom_data_size, (void) comp_formats, (void) comp_count,
            (void) input_samples_size, (void) output_samples_size;
 
        cudaStream_t stream = (cudaStream_t) vstream;
        dim3 threads_per_block(256);
        dim3 blocks((((size_x+ 7) / 8) + 255) / 256, size_y);
 
        MEASURE_KERNEL_DURATION_START(stream)
        if (size_x % 2 == 0) {
                kernel_r12l_to_rg48<uint32_t>
                    <<<blocks, threads_per_block, 0, stream>>>(
                        (uint8_t *) input_samples, (uint8_t *) output_samples,
                        size_x);
        } else {
                thread_local bool warn_print;
                if (!warn_print) {
                        fprintf(stderr,
                                "%s: Odd width %d px will use slower kernel!\n",
                                __func__, size_x);
                        warn_print = true;
                }
                kernel_r12l_to_rg48<uint16_t>
                    <<<blocks, threads_per_block, 0, stream>>>(
                        (uint8_t *) input_samples, (uint8_t *) output_samples,
                        size_x);
        }
        MEASURE_KERNEL_DURATION_STOP(stream)
 
        return cudaGetLastError();
}

References cudaGetLastError, kernel_r12l_to_rg48(), MEASURE_KERNEL_DURATION_START, and MEASURE_KERNEL_DURATION_STOP.

Here is the call graph for this function:

r12l_to_rg48_compute_blk()

template<typename store_t >

static __device__ void r12l_to_rg48_compute_blk ( const uint8_t * src, uint8_t * dst )

static

adapted variant of vc_copylineR12LtoRG48

Definition at line 344 of file kernels.cu.

{
        // load the data from in to src_u32
        auto *in_u32 = (uint32_t *) in;
        uint32_t src_u32[9];
        for (unsigned i = 0; i < sizeof src_u32 / sizeof src_u32[0]; ++i) {
                src_u32[i] = in_u32[i];
        }
 
        uint32_t dst_u32[12];
        uint8_t *dst = (uint8_t *) dst_u32;
        uint8_t *src = (uint8_t *) src_u32;
 
        // 0
        // R
        *dst++ = src[0] << 4;
        *dst++ = (src[1] << 4) | (src[0] >> 4);
        // G
        *dst++ = src[1] & 0xF0;
        *dst++ = src[2];
        // B
        *dst++ = src[3] << 4;
        *dst++ = (src[4 + 0] << 4) | (src[3] >> 4);
 
        // 1
        *dst++ = src[4 + 0] & 0xF0;
        *dst++ = src[4 + 1];
 
        *dst++ = src[4 + 2] << 4;
        *dst++ = (src[4 + 3] << 4) | (src[4 + 2] >> 4);
 
        *dst++ = src[4 + 3] & 0xF0;
        *dst++ = src[8 + 0];
 
        // 2
        *dst++ = src[8 + 1] << 4;
        *dst++ = (src[8 + 2] << 4) | (src[8 + 1] >> 4);
 
        *dst++ = src[8 + 2] & 0xF0;
        *dst++ = src[8 + 3];
 
        *dst++ = src[12 + 0] << 4;
        *dst++ = (src[12 + 1] << 4) | (src[12 + 0] >> 4);
 
        // 3
        *dst++ = src[12 + 1] & 0xF0;
        *dst++ = src[12 + 2];
 
        *dst++ = src[12 + 3] << 4;
        *dst++ = (src[16 + 0] << 4) | (src[12 + 3] >> 4);
 
        *dst++ = src[16 + 0] & 0xF0;
        *dst++ = src[16 + 1];
 
        // 4
        *dst++ = src[16 + 2] << 4;
        *dst++ = (src[16 + 3] << 4) | (src[16 + 2] >> 4);
 
        *dst++ = src[16 + 3] & 0xF0;
        *dst++ = src[20 + 0];
 
        *dst++ = src[20 + 1] << 4;
        *dst++ = (src[20 + 2] << 4) | (src[20 + 1] >> 4);
 
        // 5
        *dst++ = src[20 + 2] & 0xF0;
        *dst++ = src[20 + 3];
 
        *dst++ = src[24 + 0] << 4;
        *dst++ = (src[24 + 1] << 4) | (src[24 + 0] >> 4);
 
        *dst++ = src[24 + 1] & 0xF0;
        *dst++ = src[24 + 2];
 
        // 6
        *dst++ = src[24 + 3] << 4;
        *dst++ = (src[28 + 0] << 4) | (src[24 + 3] >> 4);
 
        *dst++ = src[28 + 0] & 0xF0;
        *dst++ = src[28 + 1];
 
        *dst++ = src[28 + 2] << 4;
        *dst++ = (src[28 + 3] << 4) | (src[28 + 2] >> 4);
 
        // 7
        *dst++ = src[28 + 3] & 0xF0;
        *dst++ = src[32 + 0];
 
        *dst++ = src[32 + 1] << 4;
        *dst++ = (src[32 + 2] << 4) | (src[32 + 1] >> 4);
 
        *dst++ = src[32 + 2] & 0xF0;
        *dst++ = src[32 + 3];
 
        // store the result
        auto *out_t = (store_t *) out;
        for (unsigned i = 0; i < sizeof dst_u32 / sizeof dst_u32[0]; ++i) {
                static_assert(sizeof(store_t) == 2 || sizeof(store_t) == 4,
                              "Just uint{16,32} supported!");
                if (sizeof(store_t) == 4) {
                        out_t[i] = dst_u32[i];
                } else {
                        out_t[2 * i] = dst_u32[i] & 0xFFFFU;
                        out_t[2 * i + 1] = dst_u32[i] >> 16;
                }
        }
}

Referenced by kernel_r12l_to_rg48().

rt48_to_r12l_compute_blk()

template<typename load_t >

static __device__ void rt48_to_r12l_compute_blk ( const uint8_t * in, uint8_t * out )

static

modified vc_copylineRG48toR12L

Definition at line 80 of file kernels.cu.

{
         // load the data from in to src_u32
         auto    *in_t = (load_t *) in;
         uint32_t src_u32[12];
         for (unsigned i = 0; i < sizeof src_u32 / sizeof src_u32[0]; ++i) {
                 static_assert(sizeof(load_t) == 2 || sizeof(load_t) == 4,
                               "Just uint{16,32} supported!");
                 if (sizeof(load_t) == 4) {
                         src_u32[i] = in_t[i];
                 } else {
                         src_u32[i] = in_t[2 * i] | in_t[2 * i + 1] << 16;
                 }
         }
 
        uint32_t dst_u32[9];
        auto *dst = (uint8_t *) dst_u32;
        auto *src = (uint8_t *) src_u32;
 
        // 0
        dst[0] = src[0] >> 4;
        dst[0] |= src[1] << 4;
        dst[1] = src[1] >> 4;
        src += 2;
 
        dst[1] |= src[0] & 0xF0;
        dst[2] = src[1];
        src += 2;
 
        dst[3] = src[0] >> 4;
        dst[3] |= src[1] << 4;
        dst[4 + 0] = src[1] >> 4;
        src += 2;
 
        // 1
        dst[4 + 0] |= src[0] & 0xF0;
        dst[4 + 1] = src[1];
        src += 2;
 
        dst[4 + 2] = src[0] >> 4;
        dst[4 + 2] |= src[1] << 4;
        dst[4 + 3] = src[1] >> 4;
        src += 2;
 
        dst[4 + 3] |= src[0] & 0xF0;
        dst[8 + 0] = src[1];
        src += 2;
 
        // 2
        dst[8 + 1] = src[0] >> 4;
        dst[8 + 1] |= src[1] << 4;
        dst[8 + 2] = src[1] >> 4;
        src += 2;
 
        dst[8 + 2] |= src[0] & 0xF0;
        dst[8 + 3] = src[1];
        src += 2;
 
        dst[12 + 0] = src[0] >> 4;
        dst[12 + 0] |= src[1] << 4;
        dst[12 + 1] = src[1] >> 4;
        src += 2;
 
        // 3
        dst[12 + 1] |= src[0] & 0xF0;
        dst[12 + 2] = src[1];
        src += 2;
 
        dst[12 + 3] = src[0] >> 4;
        dst[12 + 3] |= src[1] << 4;
        dst[16 + 0] = src[1] >> 4;
        src += 2;
 
        dst[16 + 0] |= src[0] & 0xF0;
        dst[16 + 1] = src[1];
        src += 2;
 
        // 4
        dst[16 + 2] = src[0] >> 4;
        dst[16 + 2] |= src[1] << 4;
        dst[16 + 3] = src[1] >> 4;
        src += 2;
 
        dst[16 + 3] |= src[0] & 0xF0;
        dst[20 + 0] = src[1];
        src += 2;
 
        dst[20 + 1] = src[0] >> 4;
        dst[20 + 1] |= src[1] << 4;
        dst[20 + 2] = src[1] >> 4;
        src += 2;
 
        // 5
        dst[20 + 2] |= src[0] & 0xF0;
        dst[20 + 3] = src[1];
        src += 2;
 
        dst[24 + 0] = src[0] >> 4;
        dst[24 + 0] |= src[1] << 4;
        dst[24 + 1] = src[1] >> 4;
        src += 2;
 
        dst[24 + 1] |= src[0] & 0xF0;
        dst[24 + 2] = src[1];
        src += 2;
 
        // 6
        dst[24 + 3] = src[0] >> 4;
        dst[24 + 3] |= src[1] << 4;
        dst[28 + 0] = src[1] >> 4;
        src += 2;
 
        dst[28 + 0] |= src[0] & 0xF0;
        dst[28 + 1] = src[1];
        src += 2;
 
        dst[28 + 2] = src[0] >> 4;
        dst[28 + 2] |= src[1] << 4;
        dst[28 + 3] = src[1] >> 4;
        src += 2;
 
        // 7
        dst[28 + 3] |= src[0] & 0xF0;
        dst[32 + 0] = src[1];
        src += 2;
 
        dst[32 + 1] = src[0] >> 4;
        dst[32 + 1] |= src[1] << 4;
        dst[32 + 2] = src[1] >> 4;
        src += 2;
 
        dst[32 + 2] |= src[0] & 0xF0;
        dst[32 + 3] = src[1];
        src += 2;
 
        // store the result
        auto *out_u32 = (uint32_t *) out;
        for (unsigned i = 0; i < sizeof dst_u32 / sizeof dst_u32[0]; ++i) {
                out_u32[i] = dst_u32[i];
        }
}

Referenced by kernel_rg48_to_r12l(), and rt48_to_r12l_compute_last_blk().

rt48_to_r12l_compute_last_blk()

template<typename load_t >

static __device__ void rt48_to_r12l_compute_last_blk ( uint8_t * src, uint8_t * dst, unsigned width )

static

Definition at line 224 of file kernels.cu.

{
        alignas(uint32_t) uint8_t tmp[48];
        for (unsigned i = 0; i < width * 6; ++i) {
                tmp[i] = src[i];
        }
        rt48_to_r12l_compute_blk<load_t>(tmp, dst);
}

References rt48_to_r12l_compute_blk().

Referenced by kernel_rg48_to_r12l().

Here is the call graph for this function:

Variable Documentation

log_level

volatile int log_level

extern

Definition at line 65 of file debug.cpp.

Referenced by audio_play_alsa_write_frame(), av_log_ug_callback(), vidcap_state_aja::CaptureFrames(), check_symbol_size(), configure_with(), debug_print_metadata_id(), display_caca_done(), display_caca_init(), display_decklink_put_audio_frame(), glsl_compile_link(), gpujpeg_probe_internal_codec(), keyboard_control::impl::info(), init_rtsp(), main(), module_done(), parallel_conv(), parse_opts_set_logging(), portaudio_print_help(), timed_message< log_level >::print(), print_rtsp_header(), keyboard_control::impl::read_command(), rg48_to_r12l(), keyboard_control::impl::run(), set_gamma(), signal_handler(), signal_handler(), spout_sender_register(), tfrc_done(), try_open_v4l2_device(), vidcap_file_init(), vidcap_file_worker(), vidcap_spout_init(), and video_decoder_order_output_codecs().