#include <cuda.h>
#include <iostream>
#include <string.h>
#include <stdlib.h>

/*
1. stride = 0: all threads request same value, this is where broadcasting happened
2. stride = 1: all threads request different bank
3. stride = 2: 2-way bank conflict
4. stride = 4: 4-way bank conflict
5. stride = 16: 16-way bank conflict
6. stride = 32: 32-way bank conflict

*/
__global__ void TestKernel(unsigned long long* time, int stride){
    __shared__ float shared_data[1024];
    int tid = threadIdx.x;
    unsigned long long startTime = clock();
    shared_data[tid * stride] = 4;
    for(int index = 0; index < 10; index ++){
        shared_data[tid * stride] += 4;
    }
    unsigned long long endTime = clock();
    *time = (endTime - startTime);
}


int main(int argc, const char** argv){
    /*
    Bank_Conflict -s 2
    */
    if(argc != 3){
        printf("this should be used like: ./Bank_Conflict -stride 2\n");
        return -1;
    }
    int stride = 0;
    for(int index = 0; index < argc; index++){
        if(strcmp(argv[index], "-stride") == 0){
            stride = atoi(argv[index + 1]);
        }
    }
    if(stride > 32){
        printf("share_data is an array with 1024 elements by default, so the max valid stride value is 32 by default, you are setting stride larger than 32, make sure memory access in TestKernel is valid.");
    }
    std::cout << "stride set is "<<stride<<std::endl;
    unsigned long long time;
    unsigned long long * dtime;
    cudaMalloc((void**) &dtime, sizeof(unsigned long long));
    for(int index=0; index < 10; index++){
        TestKernel<<<1, 32>>>(dtime, stride);
        cudaMemcpy(&time, dtime, sizeof(unsigned long long), cudaMemcpyDeviceToHost);
        // 14 is overhead for calling clock
        std::cout <<"Time: "<<(time - 14) / 32 << std::endl;
        std::cout << std::endl;
    } 
    cudaFree(dtime);
    cudaDeviceReset();
    return 0;
}
	.file	"tmpxft_00231824_00000000-6_bank_conflict.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3638:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3638:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z31__device_stub__Z10TestKernelPyiPyi
	.type	_Z31__device_stub__Z10TestKernelPyiPyi, @function
_Z31__device_stub__Z10TestKernelPyiPyi:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10TestKernelPyi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z31__device_stub__Z10TestKernelPyiPyi, .-_Z31__device_stub__Z10TestKernelPyiPyi
	.globl	_Z10TestKernelPyi
	.type	_Z10TestKernelPyi, @function
_Z10TestKernelPyi:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z10TestKernelPyiPyi
	.cfi_endproc
.LFE3661:
	.size	_Z10TestKernelPyi, .-_Z10TestKernelPyi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"this should be used like: ./Bank_Conflict -stride 2\n"
.LC1:
	.string	"-stride"
.LC2:
	.string	"share_data is an array with 1024 elements by default, so the max valid stride value is 32 by default, you are setting stride larger than 32, make sure memory access in TestKernel is valid."
.LC3:
	.string	"stride set is "
.LC4:
	.string	"Time: "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	cmpl	$3, %edi
	je	.L20
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	orl	$-1, %eax
	jmp	.L8
.L20:
	movq	%rsi, %r12
	xorl	%ebp, %ebp
	leaq	.LC1(%rip), %r13
	xorl	%ebx, %ebx
.L9:
	movq	(%r12,%rbp), %rdi
	movq	%r13, %rsi
	call	strcmp@PLT
	testl	%eax, %eax
	jne	.L11
	movq	8(%r12,%rbp), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
.L11:
	addq	$8, %rbp
	cmpq	$24, %rbp
	jne	.L9
	cmpl	$32, %ebx
	jle	.L12
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L12:
	leaq	_ZSt4cout(%rip), %r12
	movl	$10, %ebp
	leaq	.LC3(%rip), %rsi
	movabsq	$4294967328, %r13
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	leaq	8(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
.L14:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%r13, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	movq	8(%rsp), %rdi
	movl	%ebx, %esi
	call	_Z31__device_stub__Z10TestKernelPyiPyi
.L13:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movl	$8, %edx
	movq	%rsp, %rdi
	call	cudaMemcpy@PLT
	leaq	.LC4(%rip), %rsi
	movq	%r12, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	(%rsp), %rax
	leaq	-14(%rax), %rsi
	shrq	$5, %rsi
	call	_ZNSo9_M_insertIyEERSoT_@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	%r12, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	decl	%ebp
	jne	.L14
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaDeviceReset@PLT
	xorl	%eax, %eax
.L8:
	movq	40(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$56, %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
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC5:
	.string	"_Z10TestKernelPyi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3663:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC5(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z10TestKernelPyi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3663:
	.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:

#include <hip/hip_runtime.h>
#include <iostream>
#include <string.h>
#include <stdlib.h>

/*
1. stride = 0: all threads request same value, this is where broadcasting happened
2. stride = 1: all threads request different bank
3. stride = 2: 2-way bank conflict
4. stride = 4: 4-way bank conflict
5. stride = 16: 16-way bank conflict
6. stride = 32: 32-way bank conflict

*/
__global__ void TestKernel(unsigned long long* time, int stride){
    __shared__ float shared_data[1024];
    int tid = threadIdx.x;
    unsigned long long startTime = clock();
    shared_data[tid * stride] = 4;
    for(int index = 0; index < 10; index ++){
        shared_data[tid * stride] += 4;
    }
    unsigned long long endTime = clock();
    *time = (endTime - startTime);
}


int main(int argc, const char** argv){
    /*
    Bank_Conflict -s 2
    */
    if(argc != 3){
        printf("this should be used like: ./Bank_Conflict -stride 2\n");
        return -1;
    }
    int stride = 0;
    for(int index = 0; index < argc; index++){
        if(strcmp(argv[index], "-stride") == 0){
            stride = atoi(argv[index + 1]);
        }
    }
    if(stride > 32){
        printf("share_data is an array with 1024 elements by default, so the max valid stride value is 32 by default, you are setting stride larger than 32, make sure memory access in TestKernel is valid.");
    }
    std::cout << "stride set is "<<stride<<std::endl;
    unsigned long long time;
    unsigned long long * dtime;
    hipMalloc((void**) &dtime, sizeof(unsigned long long));
    for(int index=0; index < 10; index++){
        TestKernel<<<1, 32>>>(dtime, stride);
        hipMemcpy(&time, dtime, sizeof(unsigned long long), hipMemcpyDeviceToHost);
        // 14 is overhead for calling clock
        std::cout <<"Time: "<<(time - 14) / 32 << std::endl;
        std::cout << std::endl;
    } 
    hipFree(dtime);
    hipDeviceReset();
    return 0;
}
	.text
	.file	"bank_conflict.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z25__device_stub__TestKernelPyi # -- Begin function _Z25__device_stub__TestKernelPyi
	.type	_Z25__device_stub__TestKernelPyi,@function
_Z25__device_stub__TestKernelPyi:       # @_Z25__device_stub__TestKernelPyi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%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	$_Z10TestKernelPyi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z25__device_stub__TestKernelPyi, .Lfunc_end0-_Z25__device_stub__TestKernelPyi
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$3, %edi
	jne	.LBB1_12
# %bb.1:                                # %.preheader.preheader
	movq	%rsi, %r14
	xorl	%ebx, %ebx
	movl	$1, %r15d
.LBB1_2:                                # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movq	-8(%r14,%r15,8), %rdi
	movl	$.L.str.1, %esi
	callq	strcmp
	testl	%eax, %eax
	jne	.LBB1_4
# %bb.3:                                #   in Loop: Header=BB1_2 Depth=1
	movq	(%r14,%r15,8), %rdi
	callq	atoi
	movl	%eax, %ebx
.LBB1_4:                                #   in Loop: Header=BB1_2 Depth=1
	incq	%r15
	cmpq	$4, %r15
	jne	.LBB1_2
# %bb.5:
	cmpl	$32, %ebx
	jle	.LBB1_7
# %bb.6:
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
.LBB1_7:
	movabsq	$4294967297, %r14               # imm = 0x100000001
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$14, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %ebp
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	leaq	8(%rsp), %rdi
	movl	$8, %esi
	callq	hipMalloc
	leaq	31(%r14), %r15
	leaq	16(%rsp), %r12
.LBB1_8:                                # =>This Inner Loop Header: Depth=1
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_10
# %bb.9:                                #   in Loop: Header=BB1_8 Depth=1
	movq	8(%rsp), %rdi
	movl	%ebx, %esi
	callq	_Z25__device_stub__TestKernelPyi
.LBB1_10:                               #   in Loop: Header=BB1_8 Depth=1
	movq	8(%rsp), %rsi
	movl	$8, %edx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	16(%rsp), %rsi
	addq	$-14, %rsi
	shrq	$5, %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIyEERSoT_
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	movl	$_ZSt4cout, %eax
	addq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	decl	%ebp
	jne	.LBB1_8
# %bb.11:
	movq	8(%rsp), %rdi
	callq	hipFree
	callq	hipDeviceReset
	xorl	%eax, %eax
	jmp	.LBB1_13
.LBB1_12:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$-1, %eax
.LBB1_13:
	addq	$24, %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
	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:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10TestKernelPyi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z10TestKernelPyi,@object       # @_Z10TestKernelPyi
	.section	.rodata,"a",@progbits
	.globl	_Z10TestKernelPyi
	.p2align	3, 0x0
_Z10TestKernelPyi:
	.quad	_Z25__device_stub__TestKernelPyi
	.size	_Z10TestKernelPyi, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"-stride"
	.size	.L.str.1, 8

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"share_data is an array with 1024 elements by default, so the max valid stride value is 32 by default, you are setting stride larger than 32, make sure memory access in TestKernel is valid."
	.size	.L.str.2, 189

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"stride set is "
	.size	.L.str.3, 15

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Time: "
	.size	.L.str.4, 7

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z10TestKernelPyi"
	.size	.L__unnamed_1, 18

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"this should be used like: ./Bank_Conflict -stride 2"
	.size	.Lstr, 52

	.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 _Z25__device_stub__TestKernelPyi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10TestKernelPyi
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10TestKernelPyi       ; -- Begin function _Z10TestKernelPyi
	.globl	_Z10TestKernelPyi
	.p2align	8
	.type	_Z10TestKernelPyi,@function
_Z10TestKernelPyi:                      ; @_Z10TestKernelPyi
; %bb.0:
	s_load_b64 s[0:1], s[0:1], 0x0
	s_getreg_b32 s2, hwreg(HW_REG_SHADER_CYCLES, 0, 20)
	s_getreg_b32 s3, hwreg(HW_REG_SHADER_CYCLES, 0, 20)
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_u32 s2, s3, s2
	s_subb_u32 s3, 0, 0
	v_mov_b32_e32 v0, s2
	v_dual_mov_b32 v2, 0 :: v_dual_mov_b32 v1, s3
	s_waitcnt lgkmcnt(0)
	global_store_b64 v2, v[0:1], s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10TestKernelPyi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 12
		.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 3
		.amdhsa_next_free_sgpr 4
		.amdhsa_reserve_vcc 0
		.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	_Z10TestKernelPyi, .Lfunc_end0-_Z10TestKernelPyi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 64
; NumSgprs: 4
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 4
; NumVGPRsForWavesPerEU: 3
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 12
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10TestKernelPyi
    .private_segment_fixed_size: 0
    .sgpr_count:     4
    .sgpr_spill_count: 0
    .symbol:         _Z10TestKernelPyi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
