﻿#include <stdio.h>

#define SIZE_TEXT (sizeof(text)-1)
#define SIZE_END (sizeof(end)-1)

__device__ char text[] =
"__ bottles of beer on the wall, __ bottles of beer!\n"
"Take one down, and pass it around, ## bottles of beer on the wall!\n\n";
       
__device__ char end[] =
"01 bottle of beer on the wall, 01 bottle of beer.\n"
"Take one down and pass it around, no more bottles of beer on the wall.\n"
"\n"
"No more bottles of beer on the wall, no more bottles of beer.\n"
"Go to the store and buy some more, 99 bottles of beer on the wall.";

__global__
void bottle99(char *addr){
    int x = threadIdx.x;
	addr += x * SIZE_TEXT;

    int bottle = 99 - x;
	if (bottle == 1) {
		for (int i=0; i<SIZE_END; i++) {
			addr[i] = end[i];
		}
		addr[SIZE_END] = '\0';
	} else {
		char c1 = (bottle/10) + '0';
		char c2 = (bottle%10) + '0';

		char d1 = ((bottle-1)/10) + '0';
		char d2 = ((bottle-1)%10) + '0';

		// replace '__' by the current bottle number, and replace '##' by the next bottle number
		for (int i=0; i<SIZE_TEXT; i++) {
			int c = text[i];
			if (c == '_') {
				addr[i] = c1;
				addr[i+1] = c2;
				i++;
			} else if (c == '#') {
				addr[i] = d1;
				addr[i+1] = d2;
				i++;
			} else {
				addr[i] = text[i];
			}
        }
    }
}
      
int main()
{
    char *buffer;
    char *d_buffer;
      
    int size = SIZE_TEXT * 98 + SIZE_END + 1;

	// CPU側でメモリを確保する
    buffer = new char[size];

	// GPU側でメモリを確保する
    cudaMalloc((void**)&d_buffer, size);

	// GPU側の関数を呼び出す。（）内が、そのまま関数の引数となる
    bottle99<<<1, 99>>>(d_buffer);

	// 指定したsize分、GPUのd_bufferから、CPUのbufferへ、データを転送する
    cudaMemcpy(buffer, d_buffer, size, cudaMemcpyDeviceToHost);

	// GPU側で確保したメモリを開放する
    cudaFree(d_buffer);

	// 結果を表示する
    puts(buffer);

	// CPU側で確保したメモリを開放する
    free(buffer);

	cudaDeviceReset();
}
  
	.file	"tmpxft_0035807a_00000000-6_bottles_gpu.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z27__device_stub__Z8bottle99PcPc
	.type	_Z27__device_stub__Z8bottle99PcPc, @function
_Z27__device_stub__Z8bottle99PcPc:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z8bottle99Pc(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L2:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z27__device_stub__Z8bottle99PcPc, .-_Z27__device_stub__Z8bottle99PcPc
	.globl	_Z8bottle99Pc
	.type	_Z8bottle99Pc, @function
_Z8bottle99Pc:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z27__device_stub__Z8bottle99PcPc
	.cfi_endproc
.LFE2053:
	.size	_Z8bottle99Pc, .-_Z8bottle99Pc
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$12011, %edi
	subq	$48, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	_Znam@PLT
	leaq	8(%rsp), %rdi
	movl	$12011, %esi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967395, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movq	8(%rsp), %rdi
	call	_Z27__device_stub__Z8bottle99PcPc
.L9:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$12011, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbx, %rdi
	call	puts@PLT
	movq	%rbx, %rdi
	call	free@PLT
	call	cudaDeviceReset@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$48, %rsp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z8bottle99Pc"
.LC1:
	.string	"text"
.LC2:
	.string	"end"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.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	_Z8bottle99Pc(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movl	$121, %r9d
	leaq	_ZL4text(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$251, %r9d
	leaq	_ZL3end(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.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
.LFE2055:
	.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
	.local	_ZL3end
	.comm	_ZL3end,251,32
	.local	_ZL4text
	.comm	_ZL4text,121,32
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

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

#define SIZE_TEXT (sizeof(text)-1)
#define SIZE_END (sizeof(end)-1)

__device__ char text[] =
"__ bottles of beer on the wall, __ bottles of beer!\n"
"Take one down, and pass it around, ## bottles of beer on the wall!\n\n";
       
__device__ char end[] =
"01 bottle of beer on the wall, 01 bottle of beer.\n"
"Take one down and pass it around, no more bottles of beer on the wall.\n"
"\n"
"No more bottles of beer on the wall, no more bottles of beer.\n"
"Go to the store and buy some more, 99 bottles of beer on the wall.";

__global__
void bottle99(char *addr){
    int x = threadIdx.x;
	addr += x * SIZE_TEXT;

    int bottle = 99 - x;
	if (bottle == 1) {
		for (int i=0; i<SIZE_END; i++) {
			addr[i] = end[i];
		}
		addr[SIZE_END] = '\0';
	} else {
		char c1 = (bottle/10) + '0';
		char c2 = (bottle%10) + '0';

		char d1 = ((bottle-1)/10) + '0';
		char d2 = ((bottle-1)%10) + '0';

		// replace '__' by the current bottle number, and replace '##' by the next bottle number
		for (int i=0; i<SIZE_TEXT; i++) {
			int c = text[i];
			if (c == '_') {
				addr[i] = c1;
				addr[i+1] = c2;
				i++;
			} else if (c == '#') {
				addr[i] = d1;
				addr[i+1] = d2;
				i++;
			} else {
				addr[i] = text[i];
			}
        }
    }
}
      
int main()
{
    char *buffer;
    char *d_buffer;
      
    int size = SIZE_TEXT * 98 + SIZE_END + 1;

	// CPU側でメモリを確保する
    buffer = new char[size];

	// GPU側でメモリを確保する
    hipMalloc((void**)&d_buffer, size);

	// GPU側の関数を呼び出す。（）内が、そのまま関数の引数となる
    bottle99<<<1, 99>>>(d_buffer);

	// 指定したsize分、GPUのd_bufferから、CPUのbufferへ、データを転送する
    hipMemcpy(buffer, d_buffer, size, hipMemcpyDeviceToHost);

	// GPU側で確保したメモリを開放する
    hipFree(d_buffer);

	// 結果を表示する
    puts(buffer);

	// CPU側で確保したメモリを開放する
    free(buffer);

	hipDeviceReset();
}
  
	.text
	.file	"bottles_gpu.hip"
	.globl	_Z23__device_stub__bottle99Pc   # -- Begin function _Z23__device_stub__bottle99Pc
	.type	_Z23__device_stub__bottle99Pc,@function
_Z23__device_stub__bottle99Pc:          # @_Z23__device_stub__bottle99Pc
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%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	$_Z8bottle99Pc, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z23__device_stub__bottle99Pc, .Lfunc_end0-_Z23__device_stub__bottle99Pc
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	movl	$12011, %edi                    # imm = 0x2EEB
	callq	_Znam
	movq	%rax, %rbx
	leaq	8(%rsp), %rdi
	movl	$12011, %esi                    # imm = 0x2EEB
	callq	hipMalloc
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	98(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	8(%rsp), %rdi
	callq	_Z23__device_stub__bottle99Pc
.LBB1_2:
	movq	8(%rsp), %rsi
	movl	$12011, %edx                    # imm = 0x2EEB
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	puts
	movq	%rbx, %rdi
	callq	free
	callq	hipDeviceReset
	xorl	%eax, %eax
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.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
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	subq	$32, %rsp
	.cfi_adjust_cfa_offset 32
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8bottle99Pc, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movl	$text, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$121, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$end, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$251, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	text,@object                    # @text
	.local	text
	.comm	text,121,16
	.type	end,@object                     # @end
	.local	end
	.comm	end,251,16
	.type	_Z8bottle99Pc,@object           # @_Z8bottle99Pc
	.section	.rodata,"a",@progbits
	.globl	_Z8bottle99Pc
	.p2align	3, 0x0
_Z8bottle99Pc:
	.quad	_Z23__device_stub__bottle99Pc
	.size	_Z8bottle99Pc, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"text"
	.size	.L__unnamed_2, 5

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"end"
	.size	.L__unnamed_3, 4

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

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

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

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8bottle99Pc           ; -- Begin function _Z8bottle99Pc
	.globl	_Z8bottle99Pc
	.p2align	8
	.type	_Z8bottle99Pc,@function
_Z8bottle99Pc:                          ; @_Z8bottle99Pc
; %bb.0:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_mul_u32_u24_e32 v1, 0x78, v0
	v_mov_b32_e32 v5, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v1, s0, s0, v1
	v_add_co_ci_u32_e64 v2, null, s1, 0, s0
	s_mov_b32 s0, exec_lo
	v_cmpx_ne_u32_e32 0x62, v0
	s_xor_b32 s4, exec_lo, s0
	s_cbranch_execz .LBB0_13
; %bb.1:
	v_sub_nc_u16 v3, 0x63, v0
	v_sub_nc_u16 v0, 0x62, v0
	v_add_co_u32 v9, vcc_lo, v1, 1
	v_add_co_ci_u32_e32 v10, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_bfe_i32 v4, v3, 0, 16
	v_bfe_i32 v6, v0, 0, 16
	s_mov_b32 s3, 0
	s_mov_b64 s[0:1], 0
	s_mov_b32 s2, s3
	v_mul_i32_i24_e32 v4, 0x6667, v4
	v_mul_i32_i24_e32 v6, 0x6667, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshrrev_b32_e32 v7, 31, v4
	v_ashrrev_i32_e32 v4, 18, v4
	v_lshrrev_b32_e32 v8, 31, v6
	v_ashrrev_i32_e32 v6, 18, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u16 v4, v4, v7
	v_add_nc_u16 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u16 v7, v4, 10
	v_mul_lo_u16 v8, v6, 10
	v_add_nc_u16 v6, v6, 48
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u16 v3, v3, v7
	v_sub_nc_u16 v8, v0, v8
	v_add_nc_u16 v0, v4, 48
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u16 v7, v3, 48
	v_add_nc_u16 v8, v8, 48
.LBB0_2:                                ; %NodeBlock
                                        ; =>This Inner Loop Header: Depth=1
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, text@rel32@lo+4
	s_addc_u32 s7, s7, text@rel32@hi+12
	s_add_u32 s6, s0, s6
	s_addc_u32 s7, s1, s7
	v_add_co_u32 v3, s0, v1, s0
	global_load_u8 v11, v5, s[6:7]
	v_add_co_ci_u32_e64 v4, s0, s1, v2, s0
	s_mov_b32 s6, -1
	s_mov_b32 s0, 0
	s_mov_b32 s5, 0
                                        ; implicit-def: $sgpr1
	s_waitcnt vmcnt(0)
	v_cmp_gt_i16_e32 vcc_lo, 0x5f, v11
	s_cbranch_vccnz .LBB0_6
; %bb.3:                                ; %LeafBlock54
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_cmp_eq_u16_e32 vcc_lo, 0x5f, v11
	s_mov_b32 s5, -1
                                        ; implicit-def: $sgpr1
	s_cbranch_vccz .LBB0_5
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=1
	v_add_co_u32 v12, vcc_lo, v9, s2
	v_add_co_ci_u32_e32 v13, vcc_lo, 0, v10, vcc_lo
	global_store_b8 v[3:4], v0, off
	s_add_i32 s1, s2, 1
	s_mov_b32 s5, 0
	global_store_b8 v[12:13], v7, off
.LBB0_5:                                ; %Flow56
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s6, 0
.LBB0_6:                                ; %Flow
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s6
	s_cbranch_vccz .LBB0_8
; %bb.7:                                ; %LeafBlock
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_cmp_ne_u16_e64 s5, 35, v11
	s_mov_b32 s0, -1
                                        ; implicit-def: $sgpr1
.LBB0_8:                                ; %Flow57
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	s_and_not1_b32 vcc_lo, exec_lo, s5
	s_cbranch_vccnz .LBB0_10
; %bb.9:                                ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s0, 0
	s_mov_b32 s1, s2
	global_store_b8 v[3:4], v11, off
.LBB0_10:                               ; %Flow58
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccnz .LBB0_12
; %bb.11:                               ;   in Loop: Header=BB0_2 Depth=1
	v_add_co_u32 v11, vcc_lo, v9, s2
	v_add_co_ci_u32_e32 v12, vcc_lo, 0, v10, vcc_lo
	s_add_i32 s1, s2, 1
	s_clause 0x1
	global_store_b8 v[3:4], v6, off
	global_store_b8 v[11:12], v8, off
.LBB0_12:                               ;   in Loop: Header=BB0_2 Depth=1
	s_add_i32 s2, s1, 1
	s_cmpk_lt_u32 s1, 0x77
	s_mov_b64 s[0:1], s[2:3]
	s_cbranch_scc1 .LBB0_2
.LBB0_13:                               ; %Flow60
	s_and_not1_saveexec_b32 s0, s4
	s_cbranch_execz .LBB0_17
; %bb.14:                               ; %.preheader.preheader
	v_mov_b32_e32 v0, 0
	s_mov_b64 s[0:1], 0
.LBB0_15:                               ; %.preheader
                                        ; =>This Inner Loop Header: Depth=1
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, end@rel32@lo+4
	s_addc_u32 s3, s3, end@rel32@hi+12
	s_add_u32 s2, s0, s2
	s_addc_u32 s3, s1, s3
	v_add_co_u32 v3, vcc_lo, v1, s0
	global_load_u8 v5, v0, s[2:3]
	v_add_co_ci_u32_e32 v4, vcc_lo, s1, v2, vcc_lo
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmpk_eq_i32 s0, 0xfa
	s_waitcnt vmcnt(0)
	global_store_b8 v[3:4], v5, off
	s_cbranch_scc0 .LBB0_15
; %bb.16:
	v_mov_b32_e32 v0, 0
	global_store_b8 v[1:2], v0, off offset:250
.LBB0_17:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8bottle99Pc
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 8
		.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 14
		.amdhsa_next_free_sgpr 8
		.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	_Z8bottle99Pc, .Lfunc_end0-_Z8bottle99Pc
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 628
; NumSgprs: 10
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 10
; NumVGPRsForWavesPerEU: 14
; 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
	.protected	text                    ; @text
	.type	text,@object
	.data
	.globl	text
	.p2align	4, 0x0
text:
	.asciz	"__ bottles of beer on the wall, __ bottles of beer!\nTake one down, and pass it around, ## bottles of beer on the wall!\n\n"
	.size	text, 121

	.protected	end                     ; @end
	.type	end,@object
	.globl	end
	.p2align	4, 0x0
end:
	.asciz	"01 bottle of beer on the wall, 01 bottle of beer.\nTake one down and pass it around, no more bottles of beer on the wall.\n\nNo more bottles of beer on the wall, no more bottles of beer.\nGo to the store and buy some more, 99 bottles of beer on the wall."
	.size	end, 251

	.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 text
	.addrsig_sym end
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 8
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8bottle99Pc
    .private_segment_fixed_size: 0
    .sgpr_count:     10
    .sgpr_spill_count: 0
    .symbol:         _Z8bottle99Pc.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
