// ==========================================================================
// $Id$
// ==========================================================================
// (C)opyright: 2009-2010
//
//   Ulm University
// 
// Creator: Hendrik Lensch
// Email:   {hendrik.lensch,johannes.hanika}@uni-ulm.de
// ==========================================================================
// $Log$
// ==========================================================================

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

using namespace std;

#define MAX_BLOCKS 256
#define MAX_THREADS 128

#define RTEST // use random initialization of array
/* compute the dot product between a1 and a2. a1 and a2 are of size
 dim. The result of each thread should be stored in _dst[blockIdx.x *
 blockDim.x + threadIdx.x]. Each thread should accumulate the dot
 product of a subset of elements.
 */__global__ void dotProdKernel(float *_dst, const float* _a1, const float* _a2,
		int _dim)
{

	// program your kernel here
	float result = 0.f;
	
	const int thread_index = blockIdx.x * blockDim.x + threadIdx.x;
	
  	int elements_per_thread = (_dim + MAX_BLOCKS*MAX_THREADS -1 ) / (MAX_BLOCKS*MAX_THREADS);
	int begin = thread_index * elements_per_thread;
	int end = begin + elements_per_thread;
	
	if(end > _dim)
	  end = _dim;
	
	for(int i=begin; i<end; i++)
	{
	  result += _a1[i] * _a2[i];
	}
	
	_dst[thread_index] = result;
}

/* This program sets up two large arrays of size dim and computes the
 dot product of both arrays.

 The arrays are uploaed only once and the dot product is computed
 multiple times. While this does not make too much sense it
 demonstrated the possible speedup.  */
int main(int argc, char* argv[])
{
	// parse command line
	int acount = 1;

	if (argc < 3)
	{
		printf("usage: testDotProduct <dim> <GPU-flag [0,1]>\n");
		exit(1);
	}

	// number of elements in both vectors
	int dim = atoi(argv[acount++]);

	// flag indicating weather the CPU or the GPU version should be executed
	bool gpuVersion = atoi(argv[acount++]);

	printf("dim: %d\n", dim);

	float* cpuArray1 = new float[dim];
	float* cpuArray2 = new float[dim];

	// initialize the two arrays (either random or deterministic)
	for (int i = 0; i < dim; ++i)
	{
#ifdef RTEST 
		cpuArray1[i] = drand48();
		cpuArray2[i] = drand48();
#else 
		cpuArray1[i] = 2.0;
		cpuArray2[i] = i % 10;
#endif 
	}

	// now the gpu stuff
	float* gpuArray1;
	float* gpuArray2;
	float* gpuResult;

	float* h;

	if (gpuVersion)
	{
	const size_t input_size = dim * sizeof(float);
    	const size_t output_size = MAX_BLOCKS * MAX_THREADS * sizeof(float);
	cudaMalloc((void**) &gpuArray1, input_size);
	cudaMalloc((void**) &gpuArray2, input_size);
	cudaMalloc((void**) &gpuResult, output_size);
	cudaMemcpy(gpuArray1, cpuArray1, input_size,
			cudaMemcpyHostToDevice);
	cudaMemcpy(gpuArray2, cpuArray2, input_size,
			cudaMemcpyHostToDevice);

		// allocate an array to download the results of all threads
		h = new float[MAX_BLOCKS * MAX_THREADS];
	}

	const int num_iters = 100;
	double finalDotProduct;

	if (!gpuVersion)
	{
		printf("cpu: ");
		for (int iter = 0; iter < num_iters; ++iter)
		{
			finalDotProduct = 0.0;
			for (int i = 0; i < dim; ++i)
			{
				finalDotProduct += cpuArray1[i] * cpuArray2[i];
			}
		}
	}
	else
	{

		// Cuda version here
		printf("gpu: ");

		// a simplistic way of splitting the problem into threads
		dim3 blockGrid(MAX_BLOCKS);
		dim3 threadBlock(MAX_THREADS);

		for (int iter = 0; iter < num_iters; ++iter)
		{
			dotProdKernel<<<blockGrid, threadBlock>>>(gpuResult, gpuArray1,
					gpuArray2, dim);
			cudaThreadSynchronize();
		}

		// download and combine the results of multiple threads on the CPU

		//!!!!!!!!! missing  !!!!!!!!!!!!!!!!!!!!!!!!
	cudaMemcpy(h, gpuResult, MAX_BLOCKS * MAX_THREADS * sizeof(float), cudaMemcpyDeviceToHost);
    
    	finalDotProduct = 0.f;
    	for(int i=0; i<MAX_BLOCKS * MAX_THREADS; ++i)
      		finalDotProduct += h[i];

	}

	printf("Result: %f\n", finalDotProduct);

	if (gpuVersion)
	{

		// cleanup GPU memory

		//!!!!!!!!! missing  !!!!!!!!!!!!!!!!!!!!!!!!

		delete[] h;
	}

	delete[] cpuArray2;
	delete[] cpuArray1;

	printf("done\n");

}


	.file	"tmpxft_002ff297_00000000-6_testDotProduct.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	_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i
	.type	_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i, @function
_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i:
.LFB2052:
	.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
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%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	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%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	_Z13dotProdKernelPfPKfS1_i(%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
.LFE2052:
	.size	_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i, .-_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i
	.globl	_Z13dotProdKernelPfPKfS1_i
	.type	_Z13dotProdKernelPfPKfS1_i, @function
_Z13dotProdKernelPfPKfS1_i:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i
	.cfi_endproc
.LFE2053:
	.size	_Z13dotProdKernelPfPKfS1_i, .-_Z13dotProdKernelPfPKfS1_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"usage: testDotProduct <dim> <GPU-flag [0,1]>\n"
.LC2:
	.string	"dim: %d\n"
.LC3:
	.string	"cpu: "
.LC4:
	.string	"gpu: "
.LC5:
	.string	"Result: %f\n"
.LC6:
	.string	"done\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.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
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	cmpl	$2, %edi
	jg	.L9
	movl	$2, %edi
	leaq	.LC1(%rip), %rsi
	call	__printf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L9:
	movq	8(%rsi), %rdi
	movq	%rsi, %rbx
	call	atoi@PLT
	movq	16(%rbx), %rdi
	movl	%eax, %ebp
	call	atoi@PLT
	movl	%ebp, %edx
	movl	$2, %edi
	movslq	%ebp, %rbx
	movl	%eax, %r15d
	leaq	.LC2(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movabsq	$2305843009213693950, %rax
	cmpq	%rbx, %rax
	jb	.L10
	leaq	0(,%rbx,4), %rax
	xorl	%ebx, %ebx
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	_Znam@PLT
	movq	8(%rsp), %rdi
	movq	%rax, %r12
	call	_Znam@PLT
	movq	%rax, %r13
	jmp	.L11
.L10:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L33
	call	__cxa_throw_bad_array_new_length@PLT
.L11:
	cmpl	%ebx, %ebp
	jle	.L34
	call	drand48@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%r12,%rbx,4)
	call	drand48@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, 0(%r13,%rbx,4)
	incq	%rbx
	jmp	.L11
.L34:
	testl	%r15d, %r15d
	jne	.L14
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$100, %edx
	jmp	.L15
.L14:
	movq	8(%rsp), %rsi
	leaq	24(%rsp), %rdi
	movl	$100, %ebx
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	leaq	32(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$131072, %esi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	32(%rsp), %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$131072, %edi
	call	_Znam@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %r14
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$16777217, %eax
	movl	$1, 56(%rsp)
	salq	$8, %rax
	movl	$1, 68(%rsp)
	movq	%rax, 48(%rsp)
	addq	$-128, %rax
	movq	%rax, 60(%rsp)
	jmp	.L16
.L35:
	decl	%edx
	je	.L18
.L15:
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L19:
	cmpl	%eax, %ebp
	jle	.L35
	movss	(%r12,%rax,4), %xmm1
	mulss	0(%r13,%rax,4), %xmm1
	incq	%rax
	cvtss2sd	%xmm1, %xmm1
	addsd	%xmm1, %xmm0
	jmp	.L19
.L16:
	movl	68(%rsp), %ecx
	movq	60(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	48(%rsp), %rdi
	movl	56(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L20
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	%ebp, %ecx
	movq	40(%rsp), %rdi
	call	_Z40__device_stub__Z13dotProdKernelPfPKfS1_iPfPKfS1_i
.L20:
	call	cudaThreadSynchronize@PLT
	decl	%ebx
	jne	.L16
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movl	$131072, %edx
	movq	%r14, %rdi
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L21:
	cvtss2sd	(%r14,%rax,4), %xmm1
	incq	%rax
	addsd	%xmm1, %xmm0
	cmpq	$32768, %rax
	jne	.L21
.L18:
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	testl	%r15d, %r15d
	je	.L22
	movq	%r14, %rdi
	call	_ZdaPv@PLT
.L22:
	movq	%r13, %rdi
	call	_ZdaPv@PLT
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	xorl	%eax, %eax
	movl	$2, %edi
	leaq	.LC6(%rip), %rsi
	call	__printf_chk@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
.L33:
	call	__stack_chk_fail@PLT
.L23:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC7:
	.string	"_Z13dotProdKernelPfPKfS1_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z13dotProdKernelPfPKfS1_i(%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
.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
	.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:

// ==========================================================================
// $Id$
// ==========================================================================
// (C)opyright: 2009-2010
//
//   Ulm University
// 
// Creator: Hendrik Lensch
// Email:   {hendrik.lensch,johannes.hanika}@uni-ulm.de
// ==========================================================================
// $Log$
// ==========================================================================


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

using namespace std;

#define MAX_BLOCKS 256
#define MAX_THREADS 128

#define RTEST // use random initialization of array
/* compute the dot product between a1 and a2. a1 and a2 are of size
 dim. The result of each thread should be stored in _dst[blockIdx.x *
 blockDim.x + threadIdx.x]. Each thread should accumulate the dot
 product of a subset of elements.
 */__global__ void dotProdKernel(float *_dst, const float* _a1, const float* _a2,
		int _dim)
{

	// program your kernel here
	float result = 0.f;
	
	const int thread_index = blockIdx.x * blockDim.x + threadIdx.x;
	
  	int elements_per_thread = (_dim + MAX_BLOCKS*MAX_THREADS -1 ) / (MAX_BLOCKS*MAX_THREADS);
	int begin = thread_index * elements_per_thread;
	int end = begin + elements_per_thread;
	
	if(end > _dim)
	  end = _dim;
	
	for(int i=begin; i<end; i++)
	{
	  result += _a1[i] * _a2[i];
	}
	
	_dst[thread_index] = result;
}

/* This program sets up two large arrays of size dim and computes the
 dot product of both arrays.

 The arrays are uploaed only once and the dot product is computed
 multiple times. While this does not make too much sense it
 demonstrated the possible speedup.  */
int main(int argc, char* argv[])
{
	// parse command line
	int acount = 1;

	if (argc < 3)
	{
		printf("usage: testDotProduct <dim> <GPU-flag [0,1]>\n");
		exit(1);
	}

	// number of elements in both vectors
	int dim = atoi(argv[acount++]);

	// flag indicating weather the CPU or the GPU version should be executed
	bool gpuVersion = atoi(argv[acount++]);

	printf("dim: %d\n", dim);

	float* cpuArray1 = new float[dim];
	float* cpuArray2 = new float[dim];

	// initialize the two arrays (either random or deterministic)
	for (int i = 0; i < dim; ++i)
	{
#ifdef RTEST 
		cpuArray1[i] = drand48();
		cpuArray2[i] = drand48();
#else 
		cpuArray1[i] = 2.0;
		cpuArray2[i] = i % 10;
#endif 
	}

	// now the gpu stuff
	float* gpuArray1;
	float* gpuArray2;
	float* gpuResult;

	float* h;

	if (gpuVersion)
	{
	const size_t input_size = dim * sizeof(float);
    	const size_t output_size = MAX_BLOCKS * MAX_THREADS * sizeof(float);
	hipMalloc((void**) &gpuArray1, input_size);
	hipMalloc((void**) &gpuArray2, input_size);
	hipMalloc((void**) &gpuResult, output_size);
	hipMemcpy(gpuArray1, cpuArray1, input_size,
			hipMemcpyHostToDevice);
	hipMemcpy(gpuArray2, cpuArray2, input_size,
			hipMemcpyHostToDevice);

		// allocate an array to download the results of all threads
		h = new float[MAX_BLOCKS * MAX_THREADS];
	}

	const int num_iters = 100;
	double finalDotProduct;

	if (!gpuVersion)
	{
		printf("cpu: ");
		for (int iter = 0; iter < num_iters; ++iter)
		{
			finalDotProduct = 0.0;
			for (int i = 0; i < dim; ++i)
			{
				finalDotProduct += cpuArray1[i] * cpuArray2[i];
			}
		}
	}
	else
	{

		// Cuda version here
		printf("gpu: ");

		// a simplistic way of splitting the problem into threads
		dim3 blockGrid(MAX_BLOCKS);
		dim3 threadBlock(MAX_THREADS);

		for (int iter = 0; iter < num_iters; ++iter)
		{
			dotProdKernel<<<blockGrid, threadBlock>>>(gpuResult, gpuArray1,
					gpuArray2, dim);
			hipDeviceSynchronize();
		}

		// download and combine the results of multiple threads on the CPU

		//!!!!!!!!! missing  !!!!!!!!!!!!!!!!!!!!!!!!
	hipMemcpy(h, gpuResult, MAX_BLOCKS * MAX_THREADS * sizeof(float), hipMemcpyDeviceToHost);
    
    	finalDotProduct = 0.f;
    	for(int i=0; i<MAX_BLOCKS * MAX_THREADS; ++i)
      		finalDotProduct += h[i];

	}

	printf("Result: %f\n", finalDotProduct);

	if (gpuVersion)
	{

		// cleanup GPU memory

		//!!!!!!!!! missing  !!!!!!!!!!!!!!!!!!!!!!!!

		delete[] h;
	}

	delete[] cpuArray2;
	delete[] cpuArray1;

	printf("done\n");

}


	.text
	.file	"testDotProduct.hip"
	.globl	_Z28__device_stub__dotProdKernelPfPKfS1_i # -- Begin function _Z28__device_stub__dotProdKernelPfPKfS1_i
	.type	_Z28__device_stub__dotProdKernelPfPKfS1_i,@function
_Z28__device_stub__dotProdKernelPfPKfS1_i: # @_Z28__device_stub__dotProdKernelPfPKfS1_i
	.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	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%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	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	$_Z13dotProdKernelPfPKfS1_i, %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	_Z28__device_stub__dotProdKernelPfPKfS1_i, .Lfunc_end0-_Z28__device_stub__dotProdKernelPfPKfS1_i
	.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	$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
	cmpl	$2, %edi
	jle	.LBB1_19
# %bb.1:
	movq	%rsi, %rbx
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebp
	movq	16(%rbx), %rdi
	callq	atoi
	movl	%eax, 12(%rsp)                  # 4-byte Spill
	movl	$.L.str.1, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	movslq	%ebp, %r15
	leaq	(,%r15,4), %r12
	testl	%r15d, %r15d
	movq	$-1, %r14
	cmovnsq	%r12, %r14
	movq	%r14, %rdi
	callq	_Znam
	movq	%rax, %rbx
	movq	%r14, %rdi
	callq	_Znam
	movq	%rax, %r14
	testl	%r15d, %r15d
	jle	.LBB1_4
# %bb.2:                                # %.lr.ph.preheader
	movl	%ebp, %r13d
	xorl	%r15d, %r15d
.LBB1_3:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	drand48
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%rbx,%r15,4)
	callq	drand48
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%r14,%r15,4)
	incq	%r15
	cmpq	%r15, %r13
	jne	.LBB1_3
.LBB1_4:                                # %._crit_edge
	cmpl	$0, 12(%rsp)                    # 4-byte Folded Reload
	je	.LBB1_5
# %bb.11:
	movl	%ebp, 12(%rsp)                  # 4-byte Spill
	movabsq	$4294967424, %rbp               # imm = 0x100000080
	leaq	32(%rsp), %r13
	movq	%r13, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %r15
	movq	%r15, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$131072, %esi                   # imm = 0x20000
	callq	hipMalloc
	movq	(%r13), %rdi
	movq	%rbx, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	movq	%r14, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$131072, %edi                   # imm = 0x20000
	callq	_Znam
	movq	%rax, %r12
	movl	$.L.str.3, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$100, %r15d
	movq	%rbp, %rax
	leaq	128(%rbp), %r13
.LBB1_12:                               # =>This Inner Loop Header: Depth=1
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%rbp, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_14
# %bb.13:                               #   in Loop: Header=BB1_12 Depth=1
	movq	16(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	movl	12(%rsp), %ecx                  # 4-byte Reload
	callq	_Z28__device_stub__dotProdKernelPfPKfS1_i
.LBB1_14:                               #   in Loop: Header=BB1_12 Depth=1
	callq	hipDeviceSynchronize
	decl	%r15d
	jne	.LBB1_12
# %bb.15:
	movq	16(%rsp), %rsi
	movl	$131072, %edx                   # imm = 0x20000
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorps	%xmm0, %xmm0
	xorl	%eax, %eax
.LBB1_16:                               # =>This Inner Loop Header: Depth=1
	xorps	%xmm1, %xmm1
	cvtss2sd	(%r12,%rax,4), %xmm1
	addsd	%xmm1, %xmm0
	incq	%rax
	cmpq	$32768, %rax                    # imm = 0x8000
	jne	.LBB1_16
# %bb.17:
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
	movq	%r12, %rdi
	callq	_ZdaPv
	jmp	.LBB1_18
.LBB1_5:
	xorl	%r15d, %r15d
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	movl	%ebp, %eax
.LBB1_6:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_10 Depth 2
	xorps	%xmm0, %xmm0
	testl	%ebp, %ebp
	jle	.LBB1_7
# %bb.9:                                # %.lr.ph78.preheader
                                        #   in Loop: Header=BB1_6 Depth=1
	xorl	%ecx, %ecx
.LBB1_10:                               # %.lr.ph78
                                        #   Parent Loop BB1_6 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rbx,%rcx,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	mulss	(%r14,%rcx,4), %xmm1
	cvtss2sd	%xmm1, %xmm1
	addsd	%xmm1, %xmm0
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB1_10
.LBB1_7:                                # %._crit_edge79
                                        #   in Loop: Header=BB1_6 Depth=1
	incl	%r15d
	cmpl	$100, %r15d
	jne	.LBB1_6
# %bb.8:                                # %.thread68
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
.LBB1_18:
	movq	%r14, %rdi
	callq	_ZdaPv
	movq	%rbx, %rdi
	callq	_ZdaPv
	movl	$.Lstr, %edi
	callq	puts@PLT
	xorl	%eax, %eax
	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
	retq
.LBB1_19:
	.cfi_def_cfa_offset 96
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$1, %edi
	callq	exit
.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	$_Z13dotProdKernelPfPKfS1_i, %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	_Z13dotProdKernelPfPKfS1_i,@object # @_Z13dotProdKernelPfPKfS1_i
	.section	.rodata,"a",@progbits
	.globl	_Z13dotProdKernelPfPKfS1_i
	.p2align	3, 0x0
_Z13dotProdKernelPfPKfS1_i:
	.quad	_Z28__device_stub__dotProdKernelPfPKfS1_i
	.size	_Z13dotProdKernelPfPKfS1_i, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"dim: %d\n"
	.size	.L.str.1, 9

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"cpu: "
	.size	.L.str.2, 6

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"gpu: "
	.size	.L.str.3, 6

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Result: %f\n"
	.size	.L.str.4, 12

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z13dotProdKernelPfPKfS1_i"
	.size	.L__unnamed_1, 27

	.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	"done"
	.size	.Lstr, 5

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"usage: testDotProduct <dim> <GPU-flag [0,1]>"
	.size	.Lstr.1, 45

	.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 _Z28__device_stub__dotProdKernelPfPKfS1_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z13dotProdKernelPfPKfS1_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z13dotProdKernelPfPKfS1_i ; -- Begin function _Z13dotProdKernelPfPKfS1_i
	.globl	_Z13dotProdKernelPfPKfS1_i
	.p2align	8
	.type	_Z13dotProdKernelPfPKfS1_i,@function
_Z13dotProdKernelPfPKfS1_i:             ; @_Z13dotProdKernelPfPKfS1_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b32 s3, s[0:1], 0x18
	v_mov_b32_e32 v7, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_add_i32 s4, s3, 0x7fff
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_ashr_i32 s2, s4, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshr_b32 s2, s2, 17
	s_add_i32 s4, s4, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	s_ashr_i32 s2, s4, 15
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_mul_lo_u32 v2, v1, s2
	v_add_nc_u32_e32 v0, s2, v2
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_i32_e32 v0, s3, v0
	v_cmpx_lt_i32_e64 v2, v0
	s_cbranch_execz .LBB0_4
; %bb.1:                                ; %.lr.ph.preheader
	v_ashrrev_i32_e32 v3, 31, v2
	v_mov_b32_e32 v7, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v3, vcc_lo, s6, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v6, vcc_lo
	v_add_co_u32 v5, vcc_lo, s0, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s1, v6, vcc_lo
	s_mov_b32 s1, 0
.LBB0_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	global_load_b32 v8, v[3:4], off
	global_load_b32 v9, v[5:6], off
	v_add_nc_u32_e32 v2, 1, v2
	v_add_co_u32 v3, vcc_lo, v3, 4
	v_add_co_ci_u32_e32 v4, vcc_lo, 0, v4, vcc_lo
	v_add_co_u32 v5, vcc_lo, v5, 4
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v7, v8, v9
	v_cmp_ge_i32_e64 s0, v2, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s1, s0, s1
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB0_2
; %bb.3:                                ; %Flow
	s_or_b32 exec_lo, exec_lo, s1
.LBB0_4:                                ; %Flow42
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s2
	v_ashrrev_i32_e32 v2, 31, v1
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s5, v1, vcc_lo
	global_store_b32 v[0:1], v7, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z13dotProdKernelPfPKfS1_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 10
		.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	_Z13dotProdKernelPfPKfS1_i, .Lfunc_end0-_Z13dotProdKernelPfPKfS1_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 328
; NumSgprs: 18
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 10
; 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
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z13dotProdKernelPfPKfS1_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z13dotProdKernelPfPKfS1_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
