/*
 Copyright (c) 2016, David lu
 All rights reserved.

 Redistribution and use in source and binary forms, with or without
 modification, are permitted provided that the following conditions are met:
 * Redistributions of source code must retain the above copyright
 notice, this list of conditions and the following disclaimer.
 * Redistributions in binary form must reproduce the above copyright
 notice, this list of conditions and the following disclaimer in the
 documentation and/or other materials provided with the distribution.
 * Neither the name of the <organization> nor the
 names of its contributors may be used to endorse or promote products
 derived from this software without specific prior written permission.

 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY
 EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY
 DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */

#include <stdio.h>
#include <stdlib.h>

#define BIN_SIZE 32

using namespace std;

#define CHECK(res) if(res!=cudaSuccess){exit(-1);}

#define BLOCKNUM 1024
#define THREADNUM 128

__global__ void _k_BIT_CACU_SIGN_GPU(float_t **data_input,
		unsigned int **data_output, int num, int length) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	int out_start, in_start;

	int data_row;

	for (int j = threadid; j < num * length; j += BLOCKNUM * THREADNUM) {
		out_start = j % length;
		in_start = j % length;
		data_row = j / length;
		if (data_input[data_row][in_start] > 0)
			data_output[data_row][out_start] = 1;
	}
}

//caculate sign(I)
extern "C" void BIT_CACU_SIGN_GPU(float_t **&data, unsigned int **&out_data,
		int num, int length) {

	_k_BIT_CACU_SIGN_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data, num,
			length);

	cudaThreadSynchronize();

}

__global__ void _k_BIT_CACU_SIGN_W_GPU(float_t *data_input,
		unsigned int *data_output, int num, int length, int out_length) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	unsigned int sp[BIN_SIZE] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
			0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
	unsigned int R[BIN_SIZE] = { 1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024,
			2048, 4096, 8192, 16384, 32768, 65536, 131072, 262144, 524288,
			1048576, 2097152, 4194304, 8388608, 16777216, 33554432, 67108864,
			134217728, 268435456, 536870912, 1073741824, 2147483648 };

	int end = length - 1;

	int in_start, data;
	int data_row;
	for (int j = threadid; j < num * out_length; j += BLOCKNUM * THREADNUM) {
		data_row = j / out_length;
		in_start = j % out_length * BIN_SIZE;
		data = 0;
		for (int i = 0; i < BIN_SIZE; i++) {
			if (data_input[data_row*length + in_start + i] > 0)
				sp[i] = 1;
			else
				sp[i] = 0;
			if (in_start + i == end)
				break;
		}
		for (int i = 0; i < BIN_SIZE; i++) {
			data += R[i] * sp[i];
		}
		data_output[j] = data;
		for (int m = 0; m < BIN_SIZE; m++)
			sp[m] = 0;
	}

	//cudaFree(sp);
	//cudaFree(R);

}

//caculate sign(W)
extern "C" void BIT_CACU_SIGN_W_GPU(float_t *&data, unsigned int *&out_data,
		int w_num, int length) {

	int out_length = length / BIN_SIZE;

	_k_BIT_CACU_SIGN_W_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data, w_num,
			length, out_length);

	cudaThreadSynchronize();

}

__global__ void _k_BIT_CACU_DESIGN_GPU(float_t **data_input,
		float_t **data_output, int num, int length) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	int out_start, in_start;
	int data_row;
	for (int j = threadid; j < num * length; j += BLOCKNUM * THREADNUM) {
		data_row = j / length;
		out_start = j % length;
		in_start = j % length;
		if (abs(data_input[data_row][in_start]) < 1)
			data_output[data_row][out_start] = 1.0;
		else
			data_output[data_row][out_start] = 0.0;
	}
}

//caculate de sign(I)
extern "C" void BIT_CACU_DESIGN_GPU(float_t **&data, float_t **&out_data,
		int num, int length) {

	_k_BIT_CACU_DESIGN_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data, num,
			length);

	cudaThreadSynchronize();
}

__global__ void _k_BIT_CACU_COUNT_CONV_GPU(unsigned int *data_input,
		float_t *data_output, unsigned int *kernels, float_t *ks, float_t *a,
		int num, int kernels_num, int out_length, int motif, int block_size) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	int in_start, bit_count, index, xnor, sum;
	int data_row, data_col;

	int ks_length = out_length / kernels_num;
	int indata_length = ks_length * block_size;
	int kernel_length = block_size;

	for (int j = threadid; j < out_length * num; j += BLOCKNUM * THREADNUM) {
		data_row = j / out_length;
		data_col = j % out_length;
		in_start = data_col / kernels_num * block_size;
		index = data_col % kernels_num;
		sum = 0;
		for (int i = 0; i < block_size; i++) {
			xnor = data_input[data_row * indata_length + in_start + i]
					^ kernels[index * kernel_length + i];

			for (bit_count = 0; xnor != 0; xnor &= (xnor - 1))
				bit_count++;
			sum += bit_count;
		}
		data_output[j] = (float_t) (motif - 2 * sum)
				* ks[data_row * ks_length + data_col / kernels_num] * a[index];
	}
}

//caculate the a*ks_*(modif - 2 * bitcount(k_^x_))
//block_size is the kernel_size*kernel_size*channel / BIN_SIZE
extern "C" void BIT_CACU_COUNT_CONV_GPU(unsigned int *&data,
		unsigned int *&kernels, float_t *&ks, float_t *&a, int num, int motif,
		int kernels_num, int out_length, int block_size, float_t *&out_data) {

	_k_BIT_CACU_COUNT_CONV_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data,
			kernels, ks, a,num, kernels_num, out_length, block_size, motif);

	cudaThreadSynchronize();

}


	.file	"tmpxft_00295d97_00000000-6_bit_math.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii
	.type	_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii, @function
_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii, .-_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii
	.globl	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.type	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii, @function
_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii
	.cfi_endproc
.LFE2056:
	.size	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii, .-_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.globl	BIT_CACU_SIGN_GPU
	.type	BIT_CACU_SIGN_GPU, @function
BIT_CACU_SIGN_GPU:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	xorl	%r9d, %r9d
	movl	%ecx, %r13d
	xorl	%r8d, %r8d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%edx, %r12d
	movl	$33554433, %edx
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	salq	$7, %rdx
	movq	%rsi, %rbp
	movl	$1, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movl	$4194305, %edi
	salq	$10, %rdi
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movq	0(%rbp), %rsi
	movq	(%rbx), %rdi
	movl	%r13d, %ecx
	movl	%r12d, %edx
	call	_Z46__device_stub__Z20_k_BIT_CACU_SIGN_GPUPPfPPjiiPPfPPjii
.L9:
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	cudaThreadSynchronize@PLT
	.cfi_endproc
.LFE2027:
	.size	BIT_CACU_SIGN_GPU, .-BIT_CACU_SIGN_GPU
	.globl	_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii
	.type	_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii, @function
_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%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 168
	leaq	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	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 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L11:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii, .-_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii
	.globl	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.type	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii, @function
_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii
	.cfi_endproc
.LFE2058:
	.size	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii, .-_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.globl	BIT_CACU_SIGN_W_GPU
	.type	BIT_CACU_SIGN_W_GPU, @function
BIT_CACU_SIGN_W_GPU:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%edx, %r13d
	movl	$33554433, %edx
	xorl	%r9d, %r9d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	xorl	%r8d, %r8d
	movq	%rsi, %r12
	salq	$7, %rdx
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdi, %rbp
	movl	$4194305, %edi
	movl	$1, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	salq	$10, %rdi
	movl	%ecx, %ebx
	movl	$1, %ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L17
	movl	%ebx, %eax
	movl	$32, %ecx
	movq	(%r12), %rsi
	movq	0(%rbp), %rdi
	cltd
	idivl	%ecx
	movl	%ebx, %ecx
	movl	%r13d, %edx
	movl	%eax, %r8d
	call	_Z47__device_stub__Z22_k_BIT_CACU_SIGN_W_GPUPfPjiiiPfPjiii
.L17:
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	cudaThreadSynchronize@PLT
	.cfi_endproc
.LFE2028:
	.size	BIT_CACU_SIGN_W_GPU, .-BIT_CACU_SIGN_W_GPU
	.globl	_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii
	.type	_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii, @function
_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii:
.LFB2059:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L19:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2059:
	.size	_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii, .-_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii
	.globl	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.type	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii, @function
_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii:
.LFB2060:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii
	.cfi_endproc
.LFE2060:
	.size	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii, .-_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.globl	BIT_CACU_DESIGN_GPU
	.type	BIT_CACU_DESIGN_GPU, @function
BIT_CACU_DESIGN_GPU:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	xorl	%r9d, %r9d
	movl	%ecx, %r13d
	xorl	%r8d, %r8d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	%edx, %r12d
	movl	$33554433, %edx
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	salq	$7, %rdx
	movq	%rsi, %rbp
	movl	$1, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movl	$4194305, %edi
	salq	$10, %rdi
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	movq	0(%rbp), %rsi
	movq	(%rbx), %rdi
	movl	%r13d, %ecx
	movl	%r12d, %edx
	call	_Z48__device_stub__Z22_k_BIT_CACU_DESIGN_GPUPPfS0_iiPPfS0_ii
.L25:
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	cudaThreadSynchronize@PLT
	.cfi_endproc
.LFE2029:
	.size	BIT_CACU_DESIGN_GPU, .-BIT_CACU_DESIGN_GPU
	.globl	_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii
	.type	_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii, @function
_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii:
.LFB2061:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	248(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	64(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 240
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L27:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2061:
	.size	_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii, .-_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii
	.globl	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.type	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, @function
_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii:
.LFB2062:
	.cfi_startproc
	endbr64
	jmp	_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii
	.cfi_endproc
.LFE2062:
	.size	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, .-_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.globl	BIT_CACU_COUNT_CONV_GPU
	.type	BIT_CACU_COUNT_CONV_GPU, @function
BIT_CACU_COUNT_CONV_GPU:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%r9d, %r14d
	xorl	%r9d, %r9d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdi, %r13
	movl	$4194305, %edi
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	salq	$10, %rdi
	movq	%rcx, %r12
	movl	$1, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdx, %rbp
	movl	$33554433, %edx
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	salq	$7, %rdx
	movq	%rsi, %rbx
	movl	$1, %esi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movl	136(%rsp), %r10d
	movl	144(%rsp), %r11d
	movl	%r8d, 12(%rsp)
	xorl	%r8d, %r8d
	movq	152(%rsp), %rax
	movl	128(%rsp), %r15d
	movl	%r10d, 28(%rsp)
	movl	%r11d, 24(%rsp)
	movq	%rax, 16(%rsp)
	call	__cudaPushCallConfiguration@PLT
	movl	24(%rsp), %r11d
	movl	28(%rsp), %r10d
	testl	%eax, %eax
	jne	.L33
	pushq	%r14
	.cfi_def_cfa_offset 136
	pushq	%r11
	.cfi_def_cfa_offset 144
	pushq	%r10
	.cfi_def_cfa_offset 152
	pushq	%r15
	.cfi_def_cfa_offset 160
	movq	48(%rsp), %rax
	movl	44(%rsp), %r9d
	movq	(%r12), %r8
	movq	0(%rbp), %rcx
	movq	(%rbx), %rdx
	movq	(%rax), %rsi
	movq	0(%r13), %rdi
	call	_Z61__device_stub__Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiiiPjPfS_S0_S0_iiiii
	addq	$32, %rsp
	.cfi_def_cfa_offset 128
.L33:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	jmp	cudaThreadSynchronize@PLT
	.cfi_endproc
.LFE2030:
	.size	BIT_CACU_COUNT_CONV_GPU, .-BIT_CACU_COUNT_CONV_GPU
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii"
.LC1:
	.string	"_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii"
.LC2:
	.string	"_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii"
.LC3:
	.string	"_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2064:
	.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	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii(%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	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii(%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
.LFE2064:
	.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 (c) 2016, David lu
 All rights reserved.

 Redistribution and use in source and binary forms, with or without
 modification, are permitted provided that the following conditions are met:
 * Redistributions of source code must retain the above copyright
 notice, this list of conditions and the following disclaimer.
 * Redistributions in binary form must reproduce the above copyright
 notice, this list of conditions and the following disclaimer in the
 documentation and/or other materials provided with the distribution.
 * Neither the name of the <organization> nor the
 names of its contributors may be used to endorse or promote products
 derived from this software without specific prior written permission.

 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY
 EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY
 DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>

#define BIN_SIZE 32

using namespace std;

#define CHECK(res) if(res!=hipSuccess){exit(-1);}

#define BLOCKNUM 1024
#define THREADNUM 128

__global__ void _k_BIT_CACU_SIGN_GPU(float_t **data_input,
		unsigned int **data_output, int num, int length) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	int out_start, in_start;

	int data_row;

	for (int j = threadid; j < num * length; j += BLOCKNUM * THREADNUM) {
		out_start = j % length;
		in_start = j % length;
		data_row = j / length;
		if (data_input[data_row][in_start] > 0)
			data_output[data_row][out_start] = 1;
	}
}

//caculate sign(I)
extern "C" void BIT_CACU_SIGN_GPU(float_t **&data, unsigned int **&out_data,
		int num, int length) {

	_k_BIT_CACU_SIGN_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data, num,
			length);

	hipDeviceSynchronize();

}

__global__ void _k_BIT_CACU_SIGN_W_GPU(float_t *data_input,
		unsigned int *data_output, int num, int length, int out_length) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	unsigned int sp[BIN_SIZE] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
			0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
	unsigned int R[BIN_SIZE] = { 1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024,
			2048, 4096, 8192, 16384, 32768, 65536, 131072, 262144, 524288,
			1048576, 2097152, 4194304, 8388608, 16777216, 33554432, 67108864,
			134217728, 268435456, 536870912, 1073741824, 2147483648 };

	int end = length - 1;

	int in_start, data;
	int data_row;
	for (int j = threadid; j < num * out_length; j += BLOCKNUM * THREADNUM) {
		data_row = j / out_length;
		in_start = j % out_length * BIN_SIZE;
		data = 0;
		for (int i = 0; i < BIN_SIZE; i++) {
			if (data_input[data_row*length + in_start + i] > 0)
				sp[i] = 1;
			else
				sp[i] = 0;
			if (in_start + i == end)
				break;
		}
		for (int i = 0; i < BIN_SIZE; i++) {
			data += R[i] * sp[i];
		}
		data_output[j] = data;
		for (int m = 0; m < BIN_SIZE; m++)
			sp[m] = 0;
	}

	//cudaFree(sp);
	//cudaFree(R);

}

//caculate sign(W)
extern "C" void BIT_CACU_SIGN_W_GPU(float_t *&data, unsigned int *&out_data,
		int w_num, int length) {

	int out_length = length / BIN_SIZE;

	_k_BIT_CACU_SIGN_W_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data, w_num,
			length, out_length);

	hipDeviceSynchronize();

}

__global__ void _k_BIT_CACU_DESIGN_GPU(float_t **data_input,
		float_t **data_output, int num, int length) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	int out_start, in_start;
	int data_row;
	for (int j = threadid; j < num * length; j += BLOCKNUM * THREADNUM) {
		data_row = j / length;
		out_start = j % length;
		in_start = j % length;
		if (abs(data_input[data_row][in_start]) < 1)
			data_output[data_row][out_start] = 1.0;
		else
			data_output[data_row][out_start] = 0.0;
	}
}

//caculate de sign(I)
extern "C" void BIT_CACU_DESIGN_GPU(float_t **&data, float_t **&out_data,
		int num, int length) {

	_k_BIT_CACU_DESIGN_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data, num,
			length);

	hipDeviceSynchronize();
}

__global__ void _k_BIT_CACU_COUNT_CONV_GPU(unsigned int *data_input,
		float_t *data_output, unsigned int *kernels, float_t *ks, float_t *a,
		int num, int kernels_num, int out_length, int motif, int block_size) {

	int tid = threadIdx.x;
	int bid = blockIdx.x;

	int threadid = bid * THREADNUM + tid;

	int in_start, bit_count, index, xnor, sum;
	int data_row, data_col;

	int ks_length = out_length / kernels_num;
	int indata_length = ks_length * block_size;
	int kernel_length = block_size;

	for (int j = threadid; j < out_length * num; j += BLOCKNUM * THREADNUM) {
		data_row = j / out_length;
		data_col = j % out_length;
		in_start = data_col / kernels_num * block_size;
		index = data_col % kernels_num;
		sum = 0;
		for (int i = 0; i < block_size; i++) {
			xnor = data_input[data_row * indata_length + in_start + i]
					^ kernels[index * kernel_length + i];

			for (bit_count = 0; xnor != 0; xnor &= (xnor - 1))
				bit_count++;
			sum += bit_count;
		}
		data_output[j] = (float_t) (motif - 2 * sum)
				* ks[data_row * ks_length + data_col / kernels_num] * a[index];
	}
}

//caculate the a*ks_*(modif - 2 * bitcount(k_^x_))
//block_size is the kernel_size*kernel_size*channel / BIN_SIZE
extern "C" void BIT_CACU_COUNT_CONV_GPU(unsigned int *&data,
		unsigned int *&kernels, float_t *&ks, float_t *&a, int num, int motif,
		int kernels_num, int out_length, int block_size, float_t *&out_data) {

	_k_BIT_CACU_COUNT_CONV_GPU<<<BLOCKNUM, THREADNUM, 0>>>(data, out_data,
			kernels, ks, a,num, kernels_num, out_length, block_size, motif);

	hipDeviceSynchronize();

}


	.text
	.file	"bit_math.hip"
	.globl	_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii # -- Begin function _Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii
	.type	_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii,@function
_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii: # @_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%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	$_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii, %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	_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii, .Lfunc_end0-_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii
	.cfi_endproc
                                        # -- End function
	.globl	BIT_CACU_SIGN_GPU               # -- Begin function BIT_CACU_SIGN_GPU
	.type	BIT_CACU_SIGN_GPU,@function
BIT_CACU_SIGN_GPU:                      # @BIT_CACU_SIGN_GPU
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebx
	movl	%edx, %ebp
	movq	%rsi, %r14
	movq	%rdi, %r15
	movabsq	$4294967424, %rdx               # imm = 0x100000080
	leaq	896(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	(%r15), %rdi
	movq	(%r14), %rsi
	movl	%ebp, %edx
	movl	%ebx, %ecx
	callq	_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii
.LBB1_2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	hipDeviceSynchronize            # TAILCALL
.Lfunc_end1:
	.size	BIT_CACU_SIGN_GPU, .Lfunc_end1-BIT_CACU_SIGN_GPU
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii # -- Begin function _Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii
	.type	_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii,@function
_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii: # @_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%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	$_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii, .Lfunc_end2-_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii
	.cfi_endproc
                                        # -- End function
	.globl	BIT_CACU_SIGN_W_GPU             # -- Begin function BIT_CACU_SIGN_W_GPU
	.type	BIT_CACU_SIGN_W_GPU,@function
BIT_CACU_SIGN_W_GPU:                    # @BIT_CACU_SIGN_W_GPU
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebx
	movl	%edx, %ebp
	movq	%rsi, %r14
	movq	%rdi, %r15
	movabsq	$4294967424, %rdx               # imm = 0x100000080
	leaq	896(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	leal	31(%rbx), %r8d
	testl	%ebx, %ebx
	cmovnsl	%ebx, %r8d
	sarl	$5, %r8d
	movq	(%r15), %rdi
	movq	(%r14), %rsi
	movl	%ebp, %edx
	movl	%ebx, %ecx
	callq	_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii
.LBB3_2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	hipDeviceSynchronize            # TAILCALL
.Lfunc_end3:
	.size	BIT_CACU_SIGN_W_GPU, .Lfunc_end3-BIT_CACU_SIGN_W_GPU
	.cfi_endproc
                                        # -- End function
	.globl	_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii # -- Begin function _Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.type	_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii,@function
_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii: # @_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.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	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%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	$_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii, %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_end4:
	.size	_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii, .Lfunc_end4-_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.cfi_endproc
                                        # -- End function
	.globl	BIT_CACU_DESIGN_GPU             # -- Begin function BIT_CACU_DESIGN_GPU
	.type	BIT_CACU_DESIGN_GPU,@function
BIT_CACU_DESIGN_GPU:                    # @BIT_CACU_DESIGN_GPU
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebx
	movl	%edx, %ebp
	movq	%rsi, %r14
	movq	%rdi, %r15
	movabsq	$4294967424, %rdx               # imm = 0x100000080
	leaq	896(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_2
# %bb.1:
	movq	(%r15), %rdi
	movq	(%r14), %rsi
	movl	%ebp, %edx
	movl	%ebx, %ecx
	callq	_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii
.LBB5_2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	hipDeviceSynchronize            # TAILCALL
.Lfunc_end5:
	.size	BIT_CACU_DESIGN_GPU, .Lfunc_end5-BIT_CACU_DESIGN_GPU
	.cfi_endproc
                                        # -- End function
	.globl	_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii # -- Begin function _Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.type	_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii,@function
_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii: # @_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	224(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	232(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	240(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	248(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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_end6:
	.size	_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, .Lfunc_end6-_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.cfi_endproc
                                        # -- End function
	.globl	BIT_CACU_COUNT_CONV_GPU         # -- Begin function BIT_CACU_COUNT_CONV_GPU
	.type	BIT_CACU_COUNT_CONV_GPU,@function
BIT_CACU_COUNT_CONV_GPU:                # @BIT_CACU_COUNT_CONV_GPU
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $r9d killed $r9d def $r9
	movq	%r9, 32(%rsp)                   # 8-byte Spill
	movl	%r8d, 12(%rsp)                  # 4-byte Spill
	movq	%rcx, %r14
	movq	%rdx, %r15
	movq	%rsi, %r12
	movq	%rdi, %r13
	movq	120(%rsp), %rbp
	movl	112(%rsp), %ebx
	movl	104(%rsp), %eax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	96(%rsp), %eax
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	movabsq	$4294967424, %rdx               # imm = 0x100000080
	leaq	896(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_2
# %bb.1:
	movq	(%r13), %rdi
	movq	(%rbp), %rsi
	movq	(%r12), %rdx
	movq	(%r15), %rcx
	movq	(%r14), %r8
	movl	12(%rsp), %r9d                  # 4-byte Reload
	pushq	32(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	%rbx
	.cfi_adjust_cfa_offset 8
	pushq	32(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	48(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	callq	_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
.LBB7_2:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	hipDeviceSynchronize            # TAILCALL
.Lfunc_end7:
	.size	BIT_CACU_COUNT_CONV_GPU, .Lfunc_end7-BIT_CACU_COUNT_CONV_GPU
	.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	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii, %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	$_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %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_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__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	.LBB9_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
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii,@object # @_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.section	.rodata,"a",@progbits
	.globl	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.p2align	3, 0x0
_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii:
	.quad	_Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii
	.size	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii, 8

	.type	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii,@object # @_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.globl	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.p2align	3, 0x0
_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii:
	.quad	_Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii
	.size	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii, 8

	.type	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii,@object # @_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.globl	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.p2align	3, 0x0
_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii:
	.quad	_Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.size	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii, 8

	.type	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii,@object # @_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.globl	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.p2align	3, 0x0
_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii:
	.quad	_Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.size	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii"
	.size	.L__unnamed_2, 34

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii"
	.size	.L__unnamed_3, 35

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii"
	.size	.L__unnamed_4, 48

	.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 _Z35__device_stub___k_BIT_CACU_SIGN_GPUPPfPPjii
	.addrsig_sym _Z37__device_stub___k_BIT_CACU_SIGN_W_GPUPfPjiii
	.addrsig_sym _Z37__device_stub___k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.addrsig_sym _Z41__device_stub___k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.addrsig_sym _Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.addrsig_sym _Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.addrsig_sym _Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii ; -- Begin function _Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.globl	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
	.p2align	8
	.type	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii,@function
_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii:       ; @_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
; %bb.0:
	s_load_b64 s[4:5], s[0:1], 0x10
	v_lshl_add_u32 v0, s15, 7, v0
	s_mov_b32 s2, exec_lo
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s4, s5, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmpx_gt_i32_e64 s4, v0
	s_cbranch_execz .LBB0_5
; %bb.1:                                ; %.lr.ph.preheader
	s_ashr_i32 s6, s5, 31
	s_mov_b32 s8, 0
	s_add_i32 s2, s5, s6
	s_sub_i32 s5, 0, s5
	s_xor_b32 s7, s2, s6
	v_mov_b32_e32 v6, 1
	v_cvt_f32_u32_e32 v1, s7
	s_sub_i32 s2, 0, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, s2, v1
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mul_hi_u32 v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v1, v2
.LBB0_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v1, 31, v0
	s_mov_b32 s9, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v0, v1
	v_xor_b32_e32 v2, v2, v1
	v_xor_b32_e32 v1, s6, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v2, v5
	v_mul_lo_u32 v4, v3, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v2, v4
	v_add_nc_u32_e32 v4, 1, v3
	v_subrev_nc_u32_e32 v7, s7, v2
	v_cmp_le_u32_e32 vcc_lo, s7, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v3, v3, v4 :: v_dual_cndmask_b32 v2, v2, v7
	v_add_nc_u32_e32 v4, 1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s7, v2
	v_cndmask_b32_e32 v2, v3, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v1
	v_sub_nc_u32_e32 v3, v2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[1:2], 3, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, s0, v1
	v_add_co_ci_u32_e32 v8, vcc_lo, s1, v2, vcc_lo
	v_mad_u64_u32 v[9:10], null, s5, v3, v[0:1]
	global_load_b64 v[7:8], v[7:8], off
	v_ashrrev_i32_e32 v10, 31, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[9:10]
	s_waitcnt vmcnt(0)
	v_add_co_u32 v7, vcc_lo, v7, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v8, vcc_lo, v8, v4, vcc_lo
	global_load_b32 v7, v[7:8], off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_f32_e32 0, v7
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_add_co_u32 v1, vcc_lo, s2, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s3, v2, vcc_lo
	global_load_b64 v[1:2], v[1:2], off
	s_waitcnt vmcnt(0)
	v_add_co_u32 v1, vcc_lo, v1, v3
	v_add_co_ci_u32_e32 v2, vcc_lo, v2, v4, vcc_lo
	global_store_b32 v[1:2], v6, off
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s9
	v_add_nc_u32_e32 v0, 0x20000, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_le_i32_e32 vcc_lo, s4, v0
	s_or_b32 s8, vcc_lo, s8
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB0_2
.LBB0_5:                                ; %Flow34
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 11
		.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	_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii, .Lfunc_end0-_Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 436
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 1
; 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	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii ; -- Begin function _Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.globl	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
	.p2align	8
	.type	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii,@function
_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii:      ; @_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x10
	s_mov_b32 s8, 0
	v_lshl_add_u32 v0, s15, 7, v0
	s_mov_b32 s9, s8
	s_mov_b32 s10, s8
	s_mov_b32 s11, s8
	v_dual_mov_b32 v1, s8 :: v_dual_mov_b32 v2, s9
	v_dual_mov_b32 v3, s10 :: v_dual_mov_b32 v4, s11
	s_clause 0x7
	scratch_store_b128 off, v[1:4], off offset:128
	scratch_store_b128 off, v[1:4], off offset:112
	scratch_store_b128 off, v[1:4], off offset:96
	scratch_store_b128 off, v[1:4], off offset:80
	scratch_store_b128 off, v[1:4], off offset:64
	scratch_store_b128 off, v[1:4], off offset:48
	scratch_store_b128 off, v[1:4], off offset:32
	scratch_store_b128 off, v[1:4], off offset:16
	s_mov_b32 s2, exec_lo
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s3, s6, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cmpx_gt_i32_e64 s3, v0
	s_cbranch_execz .LBB1_9
; %bb.1:                                ; %.lr.ph.preheader
	s_ashr_i32 s4, s6, 31
	s_load_b128 s[12:15], s[0:1], 0x0
	s_add_i32 s2, s6, s4
	s_lshl_b32 s9, s6, 5
	s_xor_b32 s7, s2, s4
	v_mov_b32_e32 v6, 0
	v_cvt_f32_u32_e32 v1, s7
	s_sub_i32 s2, 0, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, s2, v1
	s_add_i32 s2, s5, -1
	v_mul_hi_u32 v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v1, v2
.LBB1_2:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_3 Depth 2
                                        ;     Child Loop BB1_5 Depth 2
                                        ;     Child Loop BB1_7 Depth 2
	v_ashrrev_i32_e32 v1, 31, v0
	s_mov_b32 s1, 0
	s_mov_b32 s0, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v0, v1
	v_xor_b32_e32 v2, v2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v2, v5
	v_mul_lo_u32 v4, v3, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v2, v4
	v_add_nc_u32_e32 v4, 1, v3
	v_subrev_nc_u32_e32 v7, s7, v2
	v_cmp_le_u32_e32 vcc_lo, s7, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v3, v3, v4 :: v_dual_cndmask_b32 v2, v2, v7
	v_xor_b32_e32 v7, s4, v1
	v_add_nc_u32_e32 v4, 1, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s7, v2
	v_cndmask_b32_e32 v2, v3, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v7
	v_sub_nc_u32_e32 v4, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_mul_lo_u32 v7, v4, s6
	v_mad_u64_u32 v[2:3], null, s9, v4, s[2:3]
	v_mul_lo_u32 v3, v4, s5
	v_lshlrev_b32_e32 v4, 5, v0
	v_sub_nc_u32_e32 v7, v0, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v2, v2, v4
	v_mov_b32_e32 v4, 16
	v_lshl_add_u32 v3, v7, 5, v3
.LBB1_3:                                ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, s0, v3
	s_add_i32 s10, s0, 1
	s_cmp_gt_u32 s0, 30
	s_cselect_b32 s11, -1, 0
	v_ashrrev_i32_e32 v8, 31, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v7, vcc_lo, s12, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v8, vcc_lo, s13, v8, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, s0, v2
	global_load_b32 v7, v[7:8], off
	s_waitcnt vmcnt(0)
	v_cmp_lt_f32_e64 s0, 0, v7
	v_cndmask_b32_e64 v7, 0, 1, s0
	s_or_b32 s0, vcc_lo, s11
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, exec_lo, s0
	s_or_b32 s1, s0, s1
	scratch_store_b32 v4, v7, off
	v_add_nc_u32_e32 v4, 4, v4
	s_mov_b32 s0, s10
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB1_3
; %bb.4:                                ; %.preheader.preheader
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	v_dual_mov_b32 v7, 16 :: v_dual_mov_b32 v2, 0
	s_mov_b64 s[0:1], 0
.LBB1_5:                                ; %.preheader
                                        ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	scratch_load_b32 v8, v7, off
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, __const._Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.R@rel32@lo+4
	s_addc_u32 s11, s11, __const._Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.R@rel32@hi+12
	s_add_u32 s10, s0, s10
	s_addc_u32 s11, s1, s11
	s_add_u32 s0, s0, 4
	s_load_b32 s10, s[10:11], 0x0
	s_addc_u32 s1, s1, 0
	s_cmpk_eq_i32 s0, 0x80
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_mad_u64_u32 v[3:4], null, v8, s10, v[2:3]
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v2, v3 :: v_dual_add_nc_u32 v7, 4, v7
	s_cbranch_scc0 .LBB1_5
; %bb.6:                                ;   in Loop: Header=BB1_2 Depth=1
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	s_mov_b32 s0, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, s14, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s15, v2, vcc_lo
	global_store_b32 v[1:2], v3, off
.LBB1_7:                                ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_add_i32 s1, s0, 16
	s_add_i32 s0, s0, 4
	scratch_store_b32 off, v6, s1
	s_cmpk_eq_i32 s0, 0x80
	s_cbranch_scc0 .LBB1_7
; %bb.8:                                ;   in Loop: Header=BB1_2 Depth=1
	v_add_nc_u32_e32 v0, 0x20000, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_le_i32_e32 vcc_lo, s3, v0
	s_or_b32 s8, vcc_lo, s8
	s_and_not1_b32 exec_lo, exec_lo, s8
	s_cbranch_execnz .LBB1_2
.LBB1_9:                                ; %Flow67
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 144
		.amdhsa_kernarg_size 28
		.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 1
		.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 9
		.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	_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii, .Lfunc_end1-_Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 720
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 144
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; 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	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii ; -- Begin function _Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.globl	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
	.p2align	8
	.type	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii,@function
_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii:     ; @_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x10
	v_lshl_add_u32 v0, s15, 7, v0
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s2, s3, s2
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmpx_gt_i32_e64 s2, v0
	s_cbranch_execz .LBB2_3
; %bb.1:                                ; %.lr.ph.preheader
	s_ashr_i32 s8, s3, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s4, s3, s8
	s_sub_i32 s3, 0, s3
	s_xor_b32 s9, s4, s8
	v_cvt_f32_u32_e32 v1, s9
	s_sub_i32 s4, 0, s9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, s4, v1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_mov_b32 s1, 0
	v_mul_hi_u32 v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v1, v1, v2
.LBB2_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v2, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v0, v2
	v_xor_b32_e32 v4, s8, v2
	v_xor_b32_e32 v2, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v2, v1
	v_mul_lo_u32 v5, v3, s9
	v_add_nc_u32_e32 v6, 1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v2, v2, v5
	v_subrev_nc_u32_e32 v5, s9, v2
	v_cmp_le_u32_e32 vcc_lo, s9, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v2, v2, v5 :: v_dual_cndmask_b32 v3, v3, v6
	v_cmp_le_u32_e32 vcc_lo, s9, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v3
	v_cndmask_b32_e32 v2, v3, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v4
	v_sub_nc_u32_e32 v2, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v3, 31, v2
	v_mad_u64_u32 v[7:8], null, s3, v2, v[0:1]
	v_add_nc_u32_e32 v0, 0x20000, v0
	v_lshlrev_b64 v[3:4], 3, v[2:3]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v8, 31, v7
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s4, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v4, vcc_lo
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	v_add_co_u32 v2, vcc_lo, s6, v3
	global_load_b64 v[5:6], v[5:6], off
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v4, vcc_lo
	s_waitcnt vmcnt(0)
	v_add_co_u32 v4, vcc_lo, v5, v7
	v_add_co_ci_u32_e32 v5, vcc_lo, v6, v8, vcc_lo
	global_load_b64 v[2:3], v[2:3], off
	global_load_b32 v4, v[4:5], off
	v_cmp_le_i32_e32 vcc_lo, s2, v0
	s_or_b32 s1, vcc_lo, s1
	s_waitcnt vmcnt(1)
	v_add_co_u32 v2, s0, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v3, s0, v3, v8, s0
	s_waitcnt vmcnt(0)
	v_cmp_lt_f32_e64 s0, |v4|, 1.0
	v_cndmask_b32_e64 v4, 0, 1.0, s0
	global_store_b32 v[2:3], v4, off
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB2_2
.LBB2_3:                                ; %Flow39
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 9
		.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_end2:
	.size	_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii, .Lfunc_end2-_Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 436
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; Occupancy: 16
; WaveLimiterHint : 1
; 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	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii ; -- Begin function _Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.globl	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
	.p2align	8
	.type	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii,@function
_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii: ; @_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
; %bb.0:
	s_load_b128 s[16:19], s[0:1], 0x28
	v_lshl_add_u32 v0, s15, 7, v0
	s_mov_b32 s2, exec_lo
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s12, s18, s16
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmpx_gt_i32_e64 s12, v0
	s_cbranch_execz .LBB3_7
; %bb.1:                                ; %.lr.ph70
	s_ashr_i32 s13, s17, 31
	s_ashr_i32 s15, s18, 31
	s_add_i32 s2, s17, s13
	s_load_b32 s21, s[0:1], 0x38
	s_xor_b32 s14, s2, s13
	s_add_i32 s2, s18, s15
	v_cvt_f32_u32_e32 v1, s14
	s_xor_b32 s16, s2, s15
	s_sub_i32 s2, 0, s14
	s_xor_b32 s4, s15, s13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s3, v1
	v_cvt_f32_u32_e32 v1, s16
	s_mul_i32 s2, s2, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_mul_hi_u32 s2, s3, s2
	s_add_i32 s20, s3, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_hi_u32 s2, s16, s20
	s_mul_i32 s3, s2, s14
	s_add_i32 s5, s2, 1
	s_sub_i32 s3, s16, s3
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_sub_i32 s6, s3, s14
	s_cmp_ge_u32 s3, s14
	s_cselect_b32 s2, s5, s2
	s_cselect_b32 s3, s6, s3
	s_add_i32 s5, s2, 1
	s_cmp_ge_u32 s3, s14
	v_cvt_u32_f32_e32 v1, v1
	s_cselect_b32 s5, s5, s2
	s_load_b64 s[2:3], s[0:1], 0x20
	s_xor_b32 s5, s5, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_sub_i32 s22, s5, s4
	s_waitcnt lgkmcnt(0)
	s_cmp_gt_i32 s21, 0
	s_cselect_b32 s23, -1, 0
	s_sub_i32 s4, 0, s16
	v_mul_lo_u32 v2, s4, v1
	s_load_b256 s[4:11], s[0:1], 0x0
	s_mov_b32 s1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_add_nc_u32_e32 v9, v1, v2
.LBB3_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB3_4 Depth 2
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v0, v1
	v_xor_b32_e32 v2, v2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v2, v9
	v_mul_lo_u32 v4, v3, s16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v2, v4
	v_add_nc_u32_e32 v4, 1, v3
	v_subrev_nc_u32_e32 v5, s16, v2
	v_cmp_le_u32_e32 vcc_lo, s16, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v3, v3, v4 :: v_dual_cndmask_b32 v2, v2, v5
	v_xor_b32_e32 v5, s15, v1
	v_add_nc_u32_e32 v4, 1, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s16, v2
	v_cndmask_b32_e32 v2, v3, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v5
	v_sub_nc_u32_e32 v3, v2, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, v3, s18
	v_sub_nc_u32_e32 v6, v0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v2, 31, v6
	v_add_nc_u32_e32 v4, v6, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v4, v4, v2
	v_xor_b32_e32 v2, s13, v2
	v_mul_hi_u32 v5, v4, s20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v7, v5, s14
	v_sub_nc_u32_e32 v4, v4, v7
	v_add_nc_u32_e32 v7, 1, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v8, s14, v4
	v_cmp_le_u32_e32 vcc_lo, s14, v4
	v_dual_cndmask_b32 v5, v5, v7 :: v_dual_cndmask_b32 v4, v4, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, 1, v5
	v_cmp_le_u32_e32 vcc_lo, s14, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v4, v5, v7, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s23
	v_xor_b32_e32 v4, v4, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v2, v4, v2
	v_mul_lo_u32 v7, v2, s17
	v_mad_u64_u32 v[4:5], null, v3, s22, v[2:3]
	s_delay_alu instid0(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v6, v7
	s_cbranch_vccnz .LBB3_5
; %bb.3:                                ; %.lr.ph65
                                        ;   in Loop: Header=BB3_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v5, v2, s21
	v_mov_b32_e32 v10, 0
	s_mov_b32 s0, s21
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[7:8], 2, v[5:6]
	v_mul_lo_u32 v5, v4, s21
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v7, vcc_lo, s8, v7
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v8, vcc_lo, s9, v8, vcc_lo
.LBB3_4:                                ; %._crit_edge
                                        ;   Parent Loop BB3_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_ashrrev_i32_e32 v6, 31, v5
	s_add_i32 s0, s0, -1
	s_cmp_eq_u32 s0, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[11:12], 2, v[5:6]
	v_add_nc_u32_e32 v5, 1, v5
	v_add_co_u32 v11, vcc_lo, s4, v11
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v12, vcc_lo, s5, v12, vcc_lo
	global_load_b32 v3, v[7:8], off
	global_load_b32 v6, v[11:12], off
	v_add_co_u32 v7, vcc_lo, v7, 4
	v_add_co_ci_u32_e32 v8, vcc_lo, 0, v8, vcc_lo
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v3, v3, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_bcnt_u32_b32 v10, v3, v10
	s_cbranch_scc0 .LBB3_4
	s_branch .LBB3_6
.LBB3_5:                                ;   in Loop: Header=BB3_2 Depth=1
	v_mov_b32_e32 v10, 0
.LBB3_6:                                ; %._crit_edge66
                                        ;   in Loop: Header=BB3_2 Depth=1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v5, 31, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, s10, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v2, vcc_lo, s2, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v3, vcc_lo
	global_load_b32 v4, v[4:5], off
	global_load_b32 v3, v[2:3], off
	v_lshlrev_b32_e32 v2, 1, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v2, s19, v2
	v_cvt_f32_i32_e32 v5, v2
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	v_add_nc_u32_e32 v0, 0x20000, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_le_i32_e32 vcc_lo, s12, v0
	v_add_co_u32 v1, s0, s6, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v2, s0, s7, v2, s0
	s_or_b32 s1, vcc_lo, s1
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v4, v4, v5
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v3, v4, v3
	global_store_b32 v[1:2], v3, off
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB3_2
.LBB3_7:                                ; %Flow111
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 60
		.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 13
		.amdhsa_next_free_sgpr 24
		.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_end3:
	.size	_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii, .Lfunc_end3-_Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 896
; NumSgprs: 26
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 13
; 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	__const._Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.R,@object ; @__const._Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.R
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
__const._Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.R:
	.long	1                               ; 0x1
	.long	2                               ; 0x2
	.long	4                               ; 0x4
	.long	8                               ; 0x8
	.long	16                              ; 0x10
	.long	32                              ; 0x20
	.long	64                              ; 0x40
	.long	128                             ; 0x80
	.long	256                             ; 0x100
	.long	512                             ; 0x200
	.long	1024                            ; 0x400
	.long	2048                            ; 0x800
	.long	4096                            ; 0x1000
	.long	8192                            ; 0x2000
	.long	16384                           ; 0x4000
	.long	32768                           ; 0x8000
	.long	65536                           ; 0x10000
	.long	131072                          ; 0x20000
	.long	262144                          ; 0x40000
	.long	524288                          ; 0x80000
	.long	1048576                         ; 0x100000
	.long	2097152                         ; 0x200000
	.long	4194304                         ; 0x400000
	.long	8388608                         ; 0x800000
	.long	16777216                        ; 0x1000000
	.long	33554432                        ; 0x2000000
	.long	67108864                        ; 0x4000000
	.long	134217728                       ; 0x8000000
	.long	268435456                       ; 0x10000000
	.long	536870912                       ; 0x20000000
	.long	1073741824                      ; 0x40000000
	.long	2147483648                      ; 0x80000000
	.size	__const._Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.R, 128

	.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:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z20_k_BIT_CACU_SIGN_GPUPPfPPjii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z20_k_BIT_CACU_SIGN_GPUPPfPPjii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .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:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 28
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii
    .private_segment_fixed_size: 144
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z22_k_BIT_CACU_SIGN_W_GPUPfPjiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .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:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z22_k_BIT_CACU_DESIGN_GPUPPfS0_ii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 60
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z26_k_BIT_CACU_COUNT_CONV_GPUPjPfS_S0_S0_iiiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
