#include <sys/time.h>
#include "cuda_runtime.h"
#include "yuyv2rgb.cuh"

static __device__ const unsigned char uchar_clipping_table[] = {
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -128 - -121
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -120 - -113
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -112 - -105
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -104 -  -97
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -96 -  -89
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -88 -  -81
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -80 -  -73
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -72 -  -65
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -64 -  -57
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -56 -  -49
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -48 -  -41
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -40 -  -33
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -32 -  -25
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -24 -  -17
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -16 -   -9
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //   -8 -   -1
	0, 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,
	183, 184, 185, 186, 187, 188, 189, 190, 191, 192, 193, 194, 195, 196, 197, 198, 199, 200, 201, 202, 203, 204, 205,
	206, 207, 208, 209, 210, 211, 212, 213, 214, 215, 216, 217, 218, 219, 220, 221, 222, 223, 224, 225, 226, 227, 228,
	229, 230, 231, 232, 233, 234, 235, 236, 237, 238, 239, 240, 241, 242, 243, 244, 245, 246, 247, 248, 249, 250, 251,
	252, 253, 254, 255, 255, 255, 255, 255, 255, 255, 255, 255, // 256-263
	255, 255, 255, 255, 255, 255, 255, 255, // 264-271
	255, 255, 255, 255, 255, 255, 255, 255, // 272-279
	255, 255, 255, 255, 255, 255, 255, 255, // 280-287
	255, 255, 255, 255, 255, 255, 255, 255, // 288-295
	255, 255, 255, 255, 255, 255, 255, 255, // 296-303
	255, 255, 255, 255, 255, 255, 255, 255, // 304-311
	255, 255, 255, 255, 255, 255, 255, 255, // 312-319
	255, 255, 255, 255, 255, 255, 255, 255, // 320-327
	255, 255, 255, 255, 255, 255, 255, 255, // 328-335
	255, 255, 255, 255, 255, 255, 255, 255, // 336-343
	255, 255, 255, 255, 255, 255, 255, 255, // 344-351
	255, 255, 255, 255, 255, 255, 255, 255, // 352-359
	255, 255, 255, 255, 255, 255, 255, 255, // 360-367
	255, 255, 255, 255, 255, 255, 255, 255, // 368-375
	255, 255, 255, 255, 255, 255, 255, 255, // 376-383
};

/** Clip a value to the range 0<val<255. For speed this is done using an
 * array, so can only cope with numbers in the range -128<val<383.
 */
static __device__ unsigned char CLIPVALUE(int val)
{
	// Old method (if)
	/* val = val < 0 ? 0 : val; */
	/* return val > 255 ? 255 : val; */
	
	// New method (array)
	const int clipping_table_offset = 128;
	return uchar_clipping_table[val + clipping_table_offset];
}

static __device__ void YUV2RGB(const unsigned char y, const unsigned char u, const unsigned char v, unsigned char* r,
                    unsigned char* g, unsigned char* b)
{
	const int y2 = (int)y;
	const int u2 = (int)u - 128;
	const int v2 = (int)v - 128;
	//std::cerr << "YUV=("<<y2<<","<<u2<<","<<v2<<")"<<std::endl;
	
	// This is the normal YUV conversion, but
	// appears to be incorrect for the firewire cameras
	/* int r2 = y2 + ( (v2*91947) >> 16); */
	/* int g2 = y2 - ( ((u2*22544) + (v2*46793)) >> 16 ); */
	/* int b2 = y2 + ( (u2*115999) >> 16); */

	// This is an adjusted version (UV spread out a bit)
	int r2 = y2 + ((v2 * 37221) >> 15);
	int g2 = y2 - (((u2 * 12975) + (v2 * 18949)) >> 15);
	int b2 = y2 + ((u2 * 66883) >> 15);
	//std::cerr << "   RGB=("<<r2<<","<<g2<<","<<b2<<")"<<std::endl;
	
	// Cap the values.
	*r = CLIPVALUE(r2);
	*g = CLIPVALUE(g2);
	*b = CLIPVALUE(b2);
}

__global__ void yuyv2rgb(char *YUV, char *RGB)
{
	unsigned char y0, y1, u, v;
	unsigned char r0, g0, b0;
	unsigned char r1, g1, b1;
	
	int nIn = blockIdx.x * blockDim.x * 4 + threadIdx.x * 4;
	y0 = (unsigned char)YUV[nIn];
	u  = (unsigned char)YUV[nIn + 1];
	y1 = (unsigned char)YUV[nIn + 2];
	v  = (unsigned char)YUV[nIn + 3];
	
	YUV2RGB(y0, u, v, &r0, &g0, &b0);
	YUV2RGB(y1, u, v, &r1, &g1, &b1);
	
	int nOut = blockIdx.x * blockDim.x * 6 + threadIdx.x * 6;
	RGB[nOut] = r0;
	RGB[nOut + 1] = g0;
	RGB[nOut + 2] = b0;
	RGB[nOut + 3] = r1;
	RGB[nOut + 4] = g1;
	RGB[nOut + 5] = b1;
}

void yuyv2rgb_cuda(char* YUV, char* RGB, int num_blocks, int block_size)
{
	yuyv2rgb<<<num_blocks, block_size>>>(YUV, RGB);
	cudaDeviceSynchronize();
}
/*
extern "C" void process_image_cuda(const void *src, int size)
{
        struct timeval ts;

        int yuv_size = size * sizeof(char);

        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tbefore copy image_data(CPU to GPU)\n", ts.tv_sec, ts.tv_usec);
        cudaError_t ret = cudaMemcpy(image_yuyv_cuda_, src, yuv_size, cudaMemcpyHostToDevice);
        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tcopy image_data(CPU to GPU) done\n", ts.tv_sec, ts.tv_usec);

        if (cudaSuccess != ret) {
                printf("cudaMemcpy fail %d\n", ret);
        }
        const int block_size = 256;
        const int num_blocks = yuv_size / (4*block_size);


        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tbefore yuyv2rgb computation\n", ts.tv_sec, ts.tv_usec);
        yuyv2rgb_cuda(image_yuyv_cuda_, image_rgb_cuda_, num_blocks, block_size);
        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tyuyv2rgb computation done\n", ts.tv_sec, ts.tv_usec);


        int rgb_size = size / 2 * 3 * sizeof(char);

        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tbefore copy image_data(GPU to CPU)\n", ts.tv_sec, ts.tv_usec);
        ret = cudaMemcpy(show_buf, image_rgb_cuda_, rgb_size, cudaMemcpyDeviceToHost);
        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tcopy image_data(GPU to CPU) done\n", ts.tv_sec, ts.tv_usec);
        printf("[%lu.%lu]\tcuda process image index = %d\n", ts.tv_sec, ts.tv_usec, ++index_pro);

        if (cudaSuccess != ret) {
                printf("cudaMemcpy fail %d\n", ret);
        }
}
*/

	.file	"tmpxft_003b7fc5_00000000-6_yuyv2rgb.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2013:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2013:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z29__device_stub__Z8yuyv2rgbPcS_PcS_
	.type	_Z29__device_stub__Z8yuyv2rgbPcS_PcS_, @function
_Z29__device_stub__Z8yuyv2rgbPcS_PcS_:
.LFB2035:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z8yuyv2rgbPcS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2035:
	.size	_Z29__device_stub__Z8yuyv2rgbPcS_PcS_, .-_Z29__device_stub__Z8yuyv2rgbPcS_PcS_
	.globl	_Z8yuyv2rgbPcS_
	.type	_Z8yuyv2rgbPcS_, @function
_Z8yuyv2rgbPcS_:
.LFB2036:
	.cfi_startproc
	endbr64
	jmp	_Z29__device_stub__Z8yuyv2rgbPcS_PcS_
	.cfi_endproc
.LFE2036:
	.size	_Z8yuyv2rgbPcS_, .-_Z8yuyv2rgbPcS_
	.globl	yuyv2rgb_cuda
	.type	yuyv2rgb_cuda, @function
yuyv2rgb_cuda:
.LFB2010:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movl	%ecx, 20(%rsp)
	movq	%rax, 24(%rsp)
	movl	28(%rsp), %ecx
	movl	%edx, 8(%rsp)
	movq	20(%rsp), %rdx
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z29__device_stub__Z8yuyv2rgbPcS_PcS_
.L9:
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE2010:
	.size	yuyv2rgb_cuda, .-yuyv2rgb_cuda
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z8yuyv2rgbPcS_"
.LC1:
	.string	"uchar_clipping_table"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2038:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z8yuyv2rgbPcS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movl	$512, %r9d
	leaq	_ZL20uchar_clipping_table(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2038:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata
	.align 32
	.type	_ZL20uchar_clipping_table, @object
	.size	_ZL20uchar_clipping_table, 512
_ZL20uchar_clipping_table:
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.ascii	"\001\002\003\004\005\006\007\b\t\n\013\f\r\016\017\020\021\022"
	.ascii	"\023\024\025\026\027\030\031\032\033\034\035\036\037 !\"#$%&"
	.ascii	"'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\]^_`a"
	.ascii	"bcdefghijklmnopqrstuvwxyz{|}~\177\200\201\202\203\204\205\206"
	.ascii	"\207\210\211\212\213\214\215\216\217\220\221\222\223\224\225"
	.ascii	"\226\227\230\231\232\233\234\235\236\237\240\241\242\243\244"
	.ascii	"\245\246\247\250\251\252\253\254\255\256\257\260\261\262\263"
	.ascii	"\264\265\266\267\270\271\272\273\274\275\276\277\300\301\302"
	.ascii	"\303\304\305\306\307\310\311\312\313\314\315\316\317\320\321"
	.ascii	"\322\323\324\325\326\327\330\331\332\333\334\335\336\337\340"
	.ascii	"\341\342\343\344\345\346\347\350\351\352\353\354\355\356\357"
	.ascii	"\360\361\362\363\364\365\366\367\370\371\372\373\374\375\376"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.ascii	"\377\377\377\377\377\377\377\377\377"
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <sys/time.h>
#include "hip/hip_runtime.h"
#include "yuyv2rgb.cuh"

static __device__ const unsigned char uchar_clipping_table[] = {
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -128 - -121
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -120 - -113
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -112 - -105
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, // -104 -  -97
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -96 -  -89
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -88 -  -81
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -80 -  -73
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -72 -  -65
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -64 -  -57
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -56 -  -49
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -48 -  -41
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -40 -  -33
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -32 -  -25
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -24 -  -17
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //  -16 -   -9
	0,
	0,
	0,
	0,
	0,
	0,
	0,
	0, //   -8 -   -1
	0, 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,
	183, 184, 185, 186, 187, 188, 189, 190, 191, 192, 193, 194, 195, 196, 197, 198, 199, 200, 201, 202, 203, 204, 205,
	206, 207, 208, 209, 210, 211, 212, 213, 214, 215, 216, 217, 218, 219, 220, 221, 222, 223, 224, 225, 226, 227, 228,
	229, 230, 231, 232, 233, 234, 235, 236, 237, 238, 239, 240, 241, 242, 243, 244, 245, 246, 247, 248, 249, 250, 251,
	252, 253, 254, 255, 255, 255, 255, 255, 255, 255, 255, 255, // 256-263
	255, 255, 255, 255, 255, 255, 255, 255, // 264-271
	255, 255, 255, 255, 255, 255, 255, 255, // 272-279
	255, 255, 255, 255, 255, 255, 255, 255, // 280-287
	255, 255, 255, 255, 255, 255, 255, 255, // 288-295
	255, 255, 255, 255, 255, 255, 255, 255, // 296-303
	255, 255, 255, 255, 255, 255, 255, 255, // 304-311
	255, 255, 255, 255, 255, 255, 255, 255, // 312-319
	255, 255, 255, 255, 255, 255, 255, 255, // 320-327
	255, 255, 255, 255, 255, 255, 255, 255, // 328-335
	255, 255, 255, 255, 255, 255, 255, 255, // 336-343
	255, 255, 255, 255, 255, 255, 255, 255, // 344-351
	255, 255, 255, 255, 255, 255, 255, 255, // 352-359
	255, 255, 255, 255, 255, 255, 255, 255, // 360-367
	255, 255, 255, 255, 255, 255, 255, 255, // 368-375
	255, 255, 255, 255, 255, 255, 255, 255, // 376-383
};

/** Clip a value to the range 0<val<255. For speed this is done using an
 * array, so can only cope with numbers in the range -128<val<383.
 */
static __device__ unsigned char CLIPVALUE(int val)
{
	// Old method (if)
	/* val = val < 0 ? 0 : val; */
	/* return val > 255 ? 255 : val; */
	
	// New method (array)
	const int clipping_table_offset = 128;
	return uchar_clipping_table[val + clipping_table_offset];
}

static __device__ void YUV2RGB(const unsigned char y, const unsigned char u, const unsigned char v, unsigned char* r,
                    unsigned char* g, unsigned char* b)
{
	const int y2 = (int)y;
	const int u2 = (int)u - 128;
	const int v2 = (int)v - 128;
	//std::cerr << "YUV=("<<y2<<","<<u2<<","<<v2<<")"<<std::endl;
	
	// This is the normal YUV conversion, but
	// appears to be incorrect for the firewire cameras
	/* int r2 = y2 + ( (v2*91947) >> 16); */
	/* int g2 = y2 - ( ((u2*22544) + (v2*46793)) >> 16 ); */
	/* int b2 = y2 + ( (u2*115999) >> 16); */

	// This is an adjusted version (UV spread out a bit)
	int r2 = y2 + ((v2 * 37221) >> 15);
	int g2 = y2 - (((u2 * 12975) + (v2 * 18949)) >> 15);
	int b2 = y2 + ((u2 * 66883) >> 15);
	//std::cerr << "   RGB=("<<r2<<","<<g2<<","<<b2<<")"<<std::endl;
	
	// Cap the values.
	*r = CLIPVALUE(r2);
	*g = CLIPVALUE(g2);
	*b = CLIPVALUE(b2);
}

__global__ void yuyv2rgb(char *YUV, char *RGB)
{
	unsigned char y0, y1, u, v;
	unsigned char r0, g0, b0;
	unsigned char r1, g1, b1;
	
	int nIn = blockIdx.x * blockDim.x * 4 + threadIdx.x * 4;
	y0 = (unsigned char)YUV[nIn];
	u  = (unsigned char)YUV[nIn + 1];
	y1 = (unsigned char)YUV[nIn + 2];
	v  = (unsigned char)YUV[nIn + 3];
	
	YUV2RGB(y0, u, v, &r0, &g0, &b0);
	YUV2RGB(y1, u, v, &r1, &g1, &b1);
	
	int nOut = blockIdx.x * blockDim.x * 6 + threadIdx.x * 6;
	RGB[nOut] = r0;
	RGB[nOut + 1] = g0;
	RGB[nOut + 2] = b0;
	RGB[nOut + 3] = r1;
	RGB[nOut + 4] = g1;
	RGB[nOut + 5] = b1;
}

void yuyv2rgb_cuda(char* YUV, char* RGB, int num_blocks, int block_size)
{
	yuyv2rgb<<<num_blocks, block_size>>>(YUV, RGB);
	hipDeviceSynchronize();
}
/*
extern "C" void process_image_cuda(const void *src, int size)
{
        struct timeval ts;

        int yuv_size = size * sizeof(char);

        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tbefore copy image_data(CPU to GPU)\n", ts.tv_sec, ts.tv_usec);
        cudaError_t ret = cudaMemcpy(image_yuyv_cuda_, src, yuv_size, cudaMemcpyHostToDevice);
        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tcopy image_data(CPU to GPU) done\n", ts.tv_sec, ts.tv_usec);

        if (cudaSuccess != ret) {
                printf("cudaMemcpy fail %d\n", ret);
        }
        const int block_size = 256;
        const int num_blocks = yuv_size / (4*block_size);


        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tbefore yuyv2rgb computation\n", ts.tv_sec, ts.tv_usec);
        yuyv2rgb_cuda(image_yuyv_cuda_, image_rgb_cuda_, num_blocks, block_size);
        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tyuyv2rgb computation done\n", ts.tv_sec, ts.tv_usec);


        int rgb_size = size / 2 * 3 * sizeof(char);

        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tbefore copy image_data(GPU to CPU)\n", ts.tv_sec, ts.tv_usec);
        ret = cudaMemcpy(show_buf, image_rgb_cuda_, rgb_size, cudaMemcpyDeviceToHost);
        gettimeofday(&ts, NULL);
        printf("[%lu.%lu]\tcopy image_data(GPU to CPU) done\n", ts.tv_sec, ts.tv_usec);
        printf("[%lu.%lu]\tcuda process image index = %d\n", ts.tv_sec, ts.tv_usec, ++index_pro);

        if (cudaSuccess != ret) {
                printf("cudaMemcpy fail %d\n", ret);
        }
}
*/

	.text
	.file	"yuyv2rgb.hip"
	.globl	_Z23__device_stub__yuyv2rgbPcS_ # -- Begin function _Z23__device_stub__yuyv2rgbPcS_
	.type	_Z23__device_stub__yuyv2rgbPcS_,@function
_Z23__device_stub__yuyv2rgbPcS_:        # @_Z23__device_stub__yuyv2rgbPcS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z8yuyv2rgbPcS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z23__device_stub__yuyv2rgbPcS_, .Lfunc_end0-_Z23__device_stub__yuyv2rgbPcS_
	.cfi_endproc
                                        # -- End function
	.globl	yuyv2rgb_cuda                   # -- Begin function yuyv2rgb_cuda
	.type	yuyv2rgb_cuda,@function
yuyv2rgb_cuda:                          # @yuyv2rgb_cuda
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsi, %rbx
	movq	%rdi, %r14
	movl	%edx, %edi
	btsq	$32, %rdi
	movl	%ecx, %edx
	btsq	$32, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	_Z23__device_stub__yuyv2rgbPcS_
.LBB1_2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	hipDeviceSynchronize            # TAILCALL
.Lfunc_end1:
	.size	yuyv2rgb_cuda, .Lfunc_end1-yuyv2rgb_cuda
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8yuyv2rgbPcS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z8yuyv2rgbPcS_,@object         # @_Z8yuyv2rgbPcS_
	.section	.rodata,"a",@progbits
	.globl	_Z8yuyv2rgbPcS_
	.p2align	3, 0x0
_Z8yuyv2rgbPcS_:
	.quad	_Z23__device_stub__yuyv2rgbPcS_
	.size	_Z8yuyv2rgbPcS_, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z8yuyv2rgbPcS_"
	.size	.L__unnamed_1, 16

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z23__device_stub__yuyv2rgbPcS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z8yuyv2rgbPcS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8yuyv2rgbPcS_         ; -- Begin function _Z8yuyv2rgbPcS_
	.globl	_Z8yuyv2rgbPcS_
	.p2align	8
	.type	_Z8yuyv2rgbPcS_,@function
_Z8yuyv2rgbPcS_:                        ; @_Z8yuyv2rgbPcS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x1c
	s_load_b128 s[0:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	v_lshlrev_b32_e32 v0, 2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v3, 31, v0
	v_add_co_u32 v2, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, _ZL20uchar_clipping_table@rel32@lo+4
	s_addc_u32 s1, s1, _ZL20uchar_clipping_table@rel32@hi+12
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL20uchar_clipping_table@rel32@lo+132
	s_addc_u32 s5, s5, _ZL20uchar_clipping_table@rel32@hi+140
	global_load_b32 v0, v[2:3], off
	s_waitcnt vmcnt(0)
	v_bfe_u32 v2, v0, 8, 8
	v_lshrrev_b32_e32 v3, 24, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, 0xffffff80, v2
	v_add_nc_u32_e32 v3, 0xffffff80, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v4, 0x32af, v2
	v_mul_i32_i24_e32 v5, 0x9165, v3
	v_mul_i32_i24_e32 v2, 0x10543, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mad_i32_i24 v3, 0x4a05, v3, v4
	v_ashrrev_i32_e32 v4, 15, v5
	s_delay_alu instid0(VALU_DEP_3)
	v_ashrrev_i32_e32 v2, 15, v2
	v_and_b32_e32 v5, 0xff, v0
	v_bfe_u32 v0, v0, 16, 8
	v_ashrrev_i32_e32 v8, 15, v3
	v_ashrrev_i32_e32 v3, 31, v4
	v_ashrrev_i32_e32 v6, 31, v2
	v_add_co_u32 v7, vcc_lo, v4, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v10, v5, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, 0, v3, vcc_lo
	v_add_co_u32 v11, vcc_lo, v2, v5
	v_add_co_ci_u32_e32 v12, vcc_lo, 0, v6, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v3, vcc_lo
	v_add_co_u32 v13, vcc_lo, v2, v0
	v_add_co_ci_u32_e32 v14, vcc_lo, 0, v6, vcc_lo
	v_add_co_u32 v2, vcc_lo, v7, s4
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v9, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, s4
	v_add_nc_u32_e32 v9, 0x80, v10
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	v_add_co_u32 v6, vcc_lo, v11, s4
	v_add_co_ci_u32_e32 v7, vcc_lo, s5, v12, vcc_lo
	s_clause 0x2
	global_load_u8 v10, v[2:3], off
	global_load_u8 v9, v9, s[0:1]
	global_load_u8 v4, v[4:5], off
	v_add_co_u32 v2, vcc_lo, v13, s4
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v14, vcc_lo
	v_sub_nc_u32_e32 v0, v0, v8
	s_clause 0x1
	global_load_u8 v5, v[6:7], off
	global_load_u8 v2, v[2:3], off
	v_add_nc_u32_e32 v0, 0x80, v0
	global_load_u8 v3, v0, s[0:1]
	v_mul_lo_u32 v0, v1, 6
	s_waitcnt vmcnt(4)
	v_lshlrev_b16 v1, 8, v9
	s_waitcnt vmcnt(3)
	v_lshlrev_b16 v4, 8, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_or_b32_e32 v1, v10, v1
	s_waitcnt vmcnt(2)
	v_or_b32_e32 v4, v5, v4
	v_ashrrev_i32_e32 v5, 31, v0
	s_waitcnt vmcnt(1)
	v_lshlrev_b16 v2, 8, v2
	v_and_b32_e32 v6, 0xffff, v1
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_lshlrev_b32_e32 v4, 16, v4
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_or_b32_e32 v2, v3, v2
	s_delay_alu instid0(VALU_DEP_3)
	v_or_b32_e32 v3, v6, v4
	s_clause 0x1
	global_store_b16 v[0:1], v2, off offset:4
	global_store_b32 v[0:1], v3, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8yuyv2rgbPcS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 15
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z8yuyv2rgbPcS_, .Lfunc_end0-_Z8yuyv2rgbPcS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 552
; NumSgprs: 18
; NumVgprs: 15
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 15
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	_ZL20uchar_clipping_table,@object ; @_ZL20uchar_clipping_table
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL20uchar_clipping_table:
	.ascii	"\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\001\002\003\004\005\006\007\b\t\n\013\f\r\016\017\020\021\022\023\024\025\026\027\030\031\032\033\034\035\036\037 !\"#$%&'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\\]^_`abcdefghijklmnopqrstuvwxyz{|}~\177\200\201\202\203\204\205\206\207\210\211\212\213\214\215\216\217\220\221\222\223\224\225\226\227\230\231\232\233\234\235\236\237\240\241\242\243\244\245\246\247\250\251\252\253\254\255\256\257\260\261\262\263\264\265\266\267\270\271\272\273\274\275\276\277\300\301\302\303\304\305\306\307\310\311\312\313\314\315\316\317\320\321\322\323\324\325\326\327\330\331\332\333\334\335\336\337\340\341\342\343\344\345\346\347\350\351\352\353\354\355\356\357\360\361\362\363\364\365\366\367\370\371\372\373\374\375\376\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377\377"
	.size	_ZL20uchar_clipping_table, 512

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8yuyv2rgbPcS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8yuyv2rgbPcS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     15
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
