/*
Copyright 2017 the arraydiff authors

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/

#include <cuda_runtime_api.h>
#include <stdint.h>
#include <stdlib.h>

__global__ void cast_u8_to_f32(
    uint32_t dim,
    const uint8_t *x,
    float *y)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx < dim) {
    y[idx] = (float)(x[idx]);
  }
}

extern "C" void arraydiff_cuda_kernel_cast_u8_to_f32(
    size_t dim,
    const uint8_t *x,
    float *y,
    cudaStream_t stream)
{
  cast_u8_to_f32<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, x, y);
}

__global__ void cast_u8x4_to_f32x4(
    uint32_t dim,
    const uint8_t *x,
    float *y)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx + 4 <= dim) {
    uint32_t i = idx >> 2;
    uchar4 vx_i = ((const uchar4 *)x)[i];
    float4 vy_i = { (float)vx_i.x, (float)vx_i.y, (float)vx_i.z, (float)vx_i.w };
    ((float4 *)y)[i] = vy_i;
  } else if (idx < dim) {
    y[idx] = (float)(x[idx]);
  }
}

extern "C" void arraydiff_cuda_kernel_cast_u8x4_to_f32x4(
    size_t dim,
    const uint8_t *x,
    float *y,
    cudaStream_t stream)
{
  cast_u8x4_to_f32x4<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, x, y);
}

	.file	"tmpxft_0032c835_00000000-6_transform.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2012:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2012:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf
	.type	_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf, @function
_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 28(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z14cast_u8_to_f32jPKhPf(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf, .-_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf
	.globl	_Z14cast_u8_to_f32jPKhPf
	.type	_Z14cast_u8_to_f32jPKhPf, @function
_Z14cast_u8_to_f32jPKhPf:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf
	.cfi_endproc
.LFE2035:
	.size	_Z14cast_u8_to_f32jPKhPf, .-_Z14cast_u8_to_f32jPKhPf
	.globl	arraydiff_cuda_kernel_cast_u8_to_f32
	.type	arraydiff_cuda_kernel_cast_u8_to_f32, @function
arraydiff_cuda_kernel_cast_u8_to_f32:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	1023(%rdi), %rax
	movq	%rdx, %r12
	movq	%rcx, %r9
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	shrq	$10, %rax
	movq	%rsi, %rbp
	movl	$4194305, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
	xorl	%r8d, %r8d
	salq	$10, %rdx
	movl	$1, %ecx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_Z38__device_stub__Z14cast_u8_to_f32jPKhPfjPKhPf
.L8:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2008:
	.size	arraydiff_cuda_kernel_cast_u8_to_f32, .-arraydiff_cuda_kernel_cast_u8_to_f32
	.globl	_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf
	.type	_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf, @function
_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf:
.LFB2036:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 28(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z18cast_u8x4_to_f32x4jPKhPf(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L11:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf, .-_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf
	.globl	_Z18cast_u8x4_to_f32x4jPKhPf
	.type	_Z18cast_u8x4_to_f32x4jPKhPf, @function
_Z18cast_u8x4_to_f32x4jPKhPf:
.LFB2037:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf
	.cfi_endproc
.LFE2037:
	.size	_Z18cast_u8x4_to_f32x4jPKhPf, .-_Z18cast_u8x4_to_f32x4jPKhPf
	.globl	arraydiff_cuda_kernel_cast_u8x4_to_f32x4
	.type	arraydiff_cuda_kernel_cast_u8x4_to_f32x4, @function
arraydiff_cuda_kernel_cast_u8x4_to_f32x4:
.LFB2009:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	1023(%rdi), %rax
	movq	%rdx, %r12
	movq	%rcx, %r9
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	shrq	$10, %rax
	movq	%rsi, %rbp
	movl	$4194305, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
	xorl	%r8d, %r8d
	salq	$10, %rdx
	movl	$1, %ecx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	addq	$32, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movl	%ebx, %edi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_Z42__device_stub__Z18cast_u8x4_to_f32x4jPKhPfjPKhPf
.L16:
	.cfi_restore_state
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2009:
	.size	arraydiff_cuda_kernel_cast_u8x4_to_f32x4, .-arraydiff_cuda_kernel_cast_u8x4_to_f32x4
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z18cast_u8x4_to_f32x4jPKhPf"
.LC1:
	.string	"_Z14cast_u8_to_f32jPKhPf"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2039:
	.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	_Z18cast_u8x4_to_f32x4jPKhPf(%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	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14cast_u8_to_f32jPKhPf(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.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
.LFE2039:
	.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
	.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:

/*
Copyright 2017 the arraydiff authors

Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at

    http://www.apache.org/licenses/LICENSE-2.0

Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.
*/


#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <stdint.h>
#include <stdlib.h>

__global__ void cast_u8_to_f32(
    uint32_t dim,
    const uint8_t *x,
    float *y)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx < dim) {
    y[idx] = (float)(x[idx]);
  }
}

extern "C" void arraydiff_cuda_kernel_cast_u8_to_f32(
    size_t dim,
    const uint8_t *x,
    float *y,
    hipStream_t stream)
{
  cast_u8_to_f32<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, x, y);
}

__global__ void cast_u8x4_to_f32x4(
    uint32_t dim,
    const uint8_t *x,
    float *y)
{
  uint32_t idx = threadIdx.x + blockDim.x * blockIdx.x;
  if (idx + 4 <= dim) {
    uint32_t i = idx >> 2;
    uchar4 vx_i = ((const uchar4 *)x)[i];
    float4 vy_i = { (float)vx_i.x, (float)vx_i.y, (float)vx_i.z, (float)vx_i.w };
    ((float4 *)y)[i] = vy_i;
  } else if (idx < dim) {
    y[idx] = (float)(x[idx]);
  }
}

extern "C" void arraydiff_cuda_kernel_cast_u8x4_to_f32x4(
    size_t dim,
    const uint8_t *x,
    float *y,
    hipStream_t stream)
{
  cast_u8x4_to_f32x4<<<(dim+1024-1)/1024, 1024, 0, stream>>>(
      dim, x, y);
}

	.text
	.file	"transform.hip"
	.globl	_Z29__device_stub__cast_u8_to_f32jPKhPf # -- Begin function _Z29__device_stub__cast_u8_to_f32jPKhPf
	.type	_Z29__device_stub__cast_u8_to_f32jPKhPf,@function
_Z29__device_stub__cast_u8_to_f32jPKhPf: # @_Z29__device_stub__cast_u8_to_f32jPKhPf
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	32(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z14cast_u8_to_f32jPKhPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z29__device_stub__cast_u8_to_f32jPKhPf, .Lfunc_end0-_Z29__device_stub__cast_u8_to_f32jPKhPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_cast_u8_to_f32 # -- Begin function arraydiff_cuda_kernel_cast_u8_to_f32
	.type	arraydiff_cuda_kernel_cast_u8_to_f32,@function
arraydiff_cuda_kernel_cast_u8_to_f32:   # @arraydiff_cuda_kernel_cast_u8_to_f32
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rcx, %r9
	movq	%rdx, %rbx
	movq	%rsi, %r14
	movq	%rdi, %r15
	leaq	1023(%rdi), %rax
	shrq	$10, %rax
	movl	%eax, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB1_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB1_2:
	.cfi_def_cfa_offset 32
	movl	%r15d, %edi
	movq	%r14, %rsi
	movq	%rbx, %rdx
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z29__device_stub__cast_u8_to_f32jPKhPf # TAILCALL
.Lfunc_end1:
	.size	arraydiff_cuda_kernel_cast_u8_to_f32, .Lfunc_end1-arraydiff_cuda_kernel_cast_u8_to_f32
	.cfi_endproc
                                        # -- End function
	.globl	_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf # -- Begin function _Z33__device_stub__cast_u8x4_to_f32x4jPKhPf
	.type	_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf,@function
_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf: # @_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	32(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z18cast_u8x4_to_f32x4jPKhPf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end2:
	.size	_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf, .Lfunc_end2-_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf
	.cfi_endproc
                                        # -- End function
	.globl	arraydiff_cuda_kernel_cast_u8x4_to_f32x4 # -- Begin function arraydiff_cuda_kernel_cast_u8x4_to_f32x4
	.type	arraydiff_cuda_kernel_cast_u8x4_to_f32x4,@function
arraydiff_cuda_kernel_cast_u8x4_to_f32x4: # @arraydiff_cuda_kernel_cast_u8x4_to_f32x4
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rcx, %r9
	movq	%rdx, %rbx
	movq	%rsi, %r14
	movq	%rdi, %r15
	leaq	1023(%rdi), %rax
	shrq	$10, %rax
	movl	%eax, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$1024, %rdx                     # imm = 0x400
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB3_2
# %bb.1:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.LBB3_2:
	.cfi_def_cfa_offset 32
	movl	%r15d, %edi
	movq	%r14, %rsi
	movq	%rbx, %rdx
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf # TAILCALL
.Lfunc_end3:
	.size	arraydiff_cuda_kernel_cast_u8x4_to_f32x4, .Lfunc_end3-arraydiff_cuda_kernel_cast_u8x4_to_f32x4
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14cast_u8_to_f32jPKhPf, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z18cast_u8x4_to_f32x4jPKhPf, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z14cast_u8_to_f32jPKhPf,@object # @_Z14cast_u8_to_f32jPKhPf
	.section	.rodata,"a",@progbits
	.globl	_Z14cast_u8_to_f32jPKhPf
	.p2align	3, 0x0
_Z14cast_u8_to_f32jPKhPf:
	.quad	_Z29__device_stub__cast_u8_to_f32jPKhPf
	.size	_Z14cast_u8_to_f32jPKhPf, 8

	.type	_Z18cast_u8x4_to_f32x4jPKhPf,@object # @_Z18cast_u8x4_to_f32x4jPKhPf
	.globl	_Z18cast_u8x4_to_f32x4jPKhPf
	.p2align	3, 0x0
_Z18cast_u8x4_to_f32x4jPKhPf:
	.quad	_Z33__device_stub__cast_u8x4_to_f32x4jPKhPf
	.size	_Z18cast_u8x4_to_f32x4jPKhPf, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z18cast_u8x4_to_f32x4jPKhPf"
	.size	.L__unnamed_2, 29

	.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 _Z29__device_stub__cast_u8_to_f32jPKhPf
	.addrsig_sym _Z33__device_stub__cast_u8x4_to_f32x4jPKhPf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z14cast_u8_to_f32jPKhPf
	.addrsig_sym _Z18cast_u8x4_to_f32x4jPKhPf
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z14cast_u8_to_f32jPKhPf ; -- Begin function _Z14cast_u8_to_f32jPKhPf
	.globl	_Z14cast_u8_to_f32jPKhPf
	.p2align	8
	.type	_Z14cast_u8_to_f32jPKhPf,@function
_Z14cast_u8_to_f32jPKhPf:               ; @_Z14cast_u8_to_f32jPKhPf
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b32 s3, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_u32_e64 s3, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x8
	v_mov_b32_e32 v2, 0
	s_waitcnt lgkmcnt(0)
	global_load_u8 v3, v1, s[0:1]
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt vmcnt(0)
	v_cvt_f32_ubyte0_e32 v2, v3
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14cast_u8_to_f32jPKhPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 4
		.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	_Z14cast_u8_to_f32jPKhPf, .Lfunc_end0-_Z14cast_u8_to_f32jPKhPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 136
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; 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
	.protected	_Z18cast_u8x4_to_f32x4jPKhPf ; -- Begin function _Z18cast_u8x4_to_f32x4jPKhPf
	.globl	_Z18cast_u8x4_to_f32x4jPKhPf
	.p2align	8
	.type	_Z18cast_u8x4_to_f32x4jPKhPf,@function
_Z18cast_u8x4_to_f32x4jPKhPf:           ; @_Z18cast_u8x4_to_f32x4jPKhPf
; %bb.0:
	s_clause 0x2
	s_load_b32 s5, s[0:1], 0x24
	s_load_b32 s4, s[0:1], 0x0
	s_load_b128 s[0:3], s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s5, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s5, v[0:1]
	s_mov_b32 s5, exec_lo
	v_add_nc_u32_e32 v0, 4, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ge_u32_e64 s4, v0
	s_xor_b32 s5, exec_lo, s5
	s_cbranch_execz .LBB1_2
; %bb.1:
	v_and_b32_e32 v0, -4, v1
	global_load_b32 v3, v0, s[0:1]
	v_lshrrev_b32_e32 v0, 2, v1
	v_mov_b32_e32 v1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 4, v[0:1]
	v_add_co_u32 v4, vcc_lo, s2, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s3, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_cvt_f32_ubyte0_e32 v0, v3
	v_cvt_f32_ubyte1_e32 v1, v3
	v_cvt_f32_ubyte2_e32 v2, v3
	v_cvt_f32_ubyte3_e32 v3, v3
	global_store_b128 v[4:5], v[0:3], off
                                        ; implicit-def: $vgpr1_vgpr2
.LBB1_2:                                ; %Flow27
	s_and_not1_saveexec_b32 s5, s5
	s_cbranch_execz .LBB1_5
; %bb.3:
	v_cmp_gt_u32_e32 vcc_lo, s4, v1
	s_and_saveexec_b32 s4, vcc_lo
	s_cbranch_execz .LBB1_5
; %bb.4:
	global_load_u8 v3, v1, s[0:1]
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_add_co_u32 v0, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	s_waitcnt vmcnt(0)
	v_cvt_f32_ubyte0_e32 v2, v3
	global_store_b32 v[0:1], v2, off
.LBB1_5:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z18cast_u8x4_to_f32x4jPKhPf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.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 6
		.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_end1:
	.size	_Z18cast_u8x4_to_f32x4jPKhPf, .Lfunc_end1-_Z18cast_u8x4_to_f32x4jPKhPf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 244
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; 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	__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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14cast_u8_to_f32jPKhPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14cast_u8_to_f32jPKhPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z18cast_u8x4_to_f32x4jPKhPf
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z18cast_u8x4_to_f32x4jPKhPf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
