/* Author: Joel Deen
* This library takes in arrays of intergers and preforms various vector operations using the GPU.
*/


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "LAVectorGPUOperations.cuh"
#include <exception>

// GPU kernel Calls ----------------------------------------------------------------------------------------------------------

// just the int add kernal, nothing special for any of the kernals
__global__ void intAdd(int* c, const int* a, const int* b, const unsigned int d)
{
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	if (i < d) {
		c[i] = a[i] + b[i];
	}
}

__global__ void intSubtract(int* c, const int* a, const int* b, const unsigned int d) {
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	if (i < d) {
		c[i] = a[i] + b[i];
	}
}

__global__ void intMultiply(int* result, const int* val1, const int val2, const unsigned int size) {
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	if (i < size) {
		result[blockIdx.x] = val1[blockIdx.x] * val2;
	}
}

__global__ void intDivide(int* result, const int* val1, const int val2, const unsigned int size) { // we are assuming that the user does some prechecking for zero
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	if (i < size) {
		result[blockIdx.x] = val1[blockIdx.x] / val2;
	}
}

// CPU Calls -----------------------------------------------------------------------------------------------------------------

int* gpuIntAdd(int* values1, int* values2, unsigned int size) {
	cudaError_t cuda_status = cudaSetDevice(0);
	if (cuda_status != cudaSuccess) {
		throw 1;
	}
	int* result = new int[size];
	int* cuda_vec1;
	int* cuda_vec2;
	int* cuda_result;

	cuda_status = cudaMalloc((void**)&cuda_vec1, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}
	cuda_status = cudaMalloc((void**)&cuda_vec2, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}
	cuda_status = cudaMalloc((void**)&cuda_result, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}

	cuda_status = cudaMemcpy(cuda_vec1, values1, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cuda_status != cudaSuccess) {
		throw 3;
	}
	cuda_status = cudaMemcpy(cuda_vec2, values2, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cuda_status != cudaSuccess) {
		throw 3;
	}
	intAdd << <1 + (size / MAX_ADDING_THREADS), MAX_ADDING_THREADS >> > (cuda_result, cuda_vec1, cuda_vec2, size);

	cuda_status = cudaGetLastError();
	if (cuda_status != cudaSuccess) {
		throw 4;
	}

	cuda_status = cudaMemcpy(result, cuda_result, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cuda_status != cudaSuccess) {
		throw 5;
	}

	cuda_status = cudaFree(cuda_vec1);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}
	cuda_status = cudaFree(cuda_vec2);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}
	cuda_status = cudaFree(cuda_result);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}

	return result;
}

int* gpuIntSubtract(int* values1, int* values2, unsigned int size) {
	cudaError_t cuda_status = cudaSetDevice(0);
	if (cuda_status != cudaSuccess) {
		throw 1;
	}
	int* result = new int[size];
	int* cuda_vec1;
	int* cuda_vec2;
	int* cuda_result;

	cuda_status = cudaMalloc((void**)&cuda_vec1, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}
	cuda_status = cudaMalloc((void**)&cuda_vec2, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}
	cuda_status = cudaMalloc((void**)&cuda_result, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}

	cuda_status = cudaMemcpy(cuda_vec1, values1, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cuda_status != cudaSuccess) {
		throw 3;
	}
	cuda_status = cudaMemcpy(cuda_vec2, values2, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cuda_status != cudaSuccess) {
		throw 3;
	}
	intSubtract << <1 + (size / MAX_ADDING_THREADS), MAX_ADDING_THREADS >> > (cuda_result, cuda_vec1, cuda_vec2, size);
	cuda_status = cudaGetLastError();
	if (cuda_status != cudaSuccess) {
		throw 4;
	}

	cuda_status = cudaMemcpy(result, cuda_result, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cuda_status != cudaSuccess) {
		throw 5;
	}

	cuda_status = cudaFree(cuda_vec1);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}
	cuda_status = cudaFree(cuda_vec2);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}
	cuda_status = cudaFree(cuda_result);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}

	return result;
}

int* gpuIntMultiply(int* values1, int value, unsigned int size) {
	cudaError_t cuda_status = cudaSetDevice(0);
	if (cuda_status != cudaSuccess) {
		throw 1;
	}
	int* result = new int[size];
	int* cuda_vec1;
	int* cuda_result;

	cuda_status = cudaMalloc((void**)&cuda_vec1, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}
	cuda_status = cudaMalloc((void**)&cuda_result, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}

	cuda_status = cudaMemcpy(cuda_vec1, values1, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cuda_status != cudaSuccess) {
		throw 3;
	}
	intMultiply << <1 + (size / MAX_ADDING_THREADS), MAX_ADDING_THREADS >> > (cuda_result, cuda_vec1, value, size);
	cuda_status = cudaGetLastError();
	if (cuda_status != cudaSuccess) {
		throw 4;
	}

	cuda_status = cudaMemcpy(result, cuda_result, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cuda_status != cudaSuccess) {
		throw 5;
	}

	cuda_status = cudaFree(cuda_vec1);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}
	cuda_status = cudaFree(cuda_result);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}

	return result;
}

int* gpuIntDivide(int* values1, int value, unsigned int size) {
	if (value == 0) {
		throw 7;
	}
	cudaError_t cuda_status = cudaSetDevice(0);
	if (cuda_status != cudaSuccess) {
		throw 1;
	}
	int* result = new int[size];
	int* cuda_vec1;
	int* cuda_result;

	cuda_status = cudaMalloc((void**)&cuda_vec1, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}
	cuda_status = cudaMalloc((void**)&cuda_result, size * sizeof(int));
	if (cuda_status != cudaSuccess) {
		throw 2;
	}

	cuda_status = cudaMemcpy(cuda_vec1, values1, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cuda_status != cudaSuccess) {
		throw 3;
	}
	intMultiply << <1 + (size / MAX_ADDING_THREADS), MAX_ADDING_THREADS >> > (cuda_result, cuda_vec1, value, size);
	cuda_status = cudaGetLastError();
	if (cuda_status != cudaSuccess) {
		throw 4;
	}

	cuda_status = cudaMemcpy(result, cuda_result, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cuda_status != cudaSuccess) {
		throw 5;
	}

	cuda_status = cudaFree(cuda_vec1);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}
	cuda_status = cudaFree(cuda_result);
	if (cuda_status != cudaSuccess) {
		throw 6;
	}

	return result;
}
	.file	"tmpxft_00259649_00000000-6_LAVectorGPUOperations.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2078:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2078:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j
	.type	_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j, @function
_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j:
.LFB2100:
	.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	_Z6intAddPiPKiS1_j(%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
.LFE2100:
	.size	_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j, .-_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j
	.globl	_Z6intAddPiPKiS1_j
	.type	_Z6intAddPiPKiS1_j, @function
_Z6intAddPiPKiS1_j:
.LFB2101:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j
	.cfi_endproc
.LFE2101:
	.size	_Z6intAddPiPKiS1_j, .-_Z6intAddPiPKiS1_j
	.globl	_Z9gpuIntAddPiS_j
	.type	_Z9gpuIntAddPiS_j, @function
_Z9gpuIntAddPiS_j:
.LFB2072:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %r14
	xorl	%edi, %edi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	call	cudaSetDevice@PLT
	testl	%eax, %eax
	je	.L9
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$1, (%rax)
	movq	%rax, %rdi
	jmp	.L37
.L9:
	movl	%ebp, %ebx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	_Znam@PLT
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%rax, %r12
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L39
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L39
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L15
.L39:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$2, (%rax)
	movq	%rax, %rdi
	jmp	.L37
.L15:
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r14, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	jne	.L40
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r13, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L19
.L40:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$3, (%rax)
	movq	%rax, %rdi
	jmp	.L37
.L19:
	movl	%ebp, %eax
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	shrl	$10, %eax
	salq	$10, %rdx
	movl	$1, %ecx
	incl	%eax
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	%ebp, %ecx
	movq	24(%rsp), %rdi
	call	_Z32__device_stub__Z6intAddPiPKiS1_jPiPKiS1_j
.L21:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L22
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$4, (%rax)
	movq	%rax, %rdi
	jmp	.L37
.L22:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L24
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$5, (%rax)
	movq	%rax, %rdi
	jmp	.L37
.L24:
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	jne	.L36
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	jne	.L36
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	je	.L8
.L36:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$6, (%rax)
	movq	%rax, %rdi
.L37:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L35
	xorl	%edx, %edx
	leaq	_ZTIi(%rip), %rsi
	call	__cxa_throw@PLT
.L8:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
.L35:
	call	__stack_chk_fail@PLT
.L32:
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z9gpuIntAddPiS_j, .-_Z9gpuIntAddPiS_j
	.globl	_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j
	.type	_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j, @function
_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j:
.LFB2102:
	.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	.L41
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11intSubtractPiPKiS1_j(%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
.L41:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L43
	call	__stack_chk_fail@PLT
.L43:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2102:
	.size	_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j, .-_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j
	.globl	_Z11intSubtractPiPKiS1_j
	.type	_Z11intSubtractPiPKiS1_j, @function
_Z11intSubtractPiPKiS1_j:
.LFB2103:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j
	.cfi_endproc
.LFE2103:
	.size	_Z11intSubtractPiPKiS1_j, .-_Z11intSubtractPiPKiS1_j
	.globl	_Z14gpuIntSubtractPiS_j
	.type	_Z14gpuIntSubtractPiS_j, @function
_Z14gpuIntSubtractPiS_j:
.LFB2073:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %r14
	xorl	%edi, %edi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	call	cudaSetDevice@PLT
	testl	%eax, %eax
	je	.L47
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$1, (%rax)
	movq	%rax, %rdi
	jmp	.L75
.L47:
	movl	%ebp, %ebx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	_Znam@PLT
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%rax, %r12
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L77
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L77
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L53
.L77:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$2, (%rax)
	movq	%rax, %rdi
	jmp	.L75
.L53:
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r14, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	jne	.L78
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r13, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L57
.L78:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$3, (%rax)
	movq	%rax, %rdi
	jmp	.L75
.L57:
	movl	%ebp, %eax
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	shrl	$10, %eax
	salq	$10, %rdx
	movl	$1, %ecx
	incl	%eax
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L59
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	%ebp, %ecx
	movq	24(%rsp), %rdi
	call	_Z38__device_stub__Z11intSubtractPiPKiS1_jPiPKiS1_j
.L59:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L60
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$4, (%rax)
	movq	%rax, %rdi
	jmp	.L75
.L60:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L62
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$5, (%rax)
	movq	%rax, %rdi
	jmp	.L75
.L62:
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	jne	.L74
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	jne	.L74
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	je	.L46
.L74:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$6, (%rax)
	movq	%rax, %rdi
.L75:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L73
	xorl	%edx, %edx
	leaq	_ZTIi(%rip), %rsi
	call	__cxa_throw@PLT
.L46:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L70
.L73:
	call	__stack_chk_fail@PLT
.L70:
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2073:
	.size	_Z14gpuIntSubtractPiS_j, .-_Z14gpuIntSubtractPiS_j
	.globl	_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij
	.type	_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij, @function
_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij:
.LFB2104:
	.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	.L79
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z11intMultiplyPiPKiij(%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
.L79:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L81
	call	__stack_chk_fail@PLT
.L81:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2104:
	.size	_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij, .-_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij
	.globl	_Z11intMultiplyPiPKiij
	.type	_Z11intMultiplyPiPKiij, @function
_Z11intMultiplyPiPKiij:
.LFB2105:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij
	.cfi_endproc
.LFE2105:
	.size	_Z11intMultiplyPiPKiij, .-_Z11intMultiplyPiPKiij
	.globl	_Z14gpuIntMultiplyPiij
	.type	_Z14gpuIntMultiplyPiij, @function
_Z14gpuIntMultiplyPiij:
.LFB2074:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdi, %r14
	xorl	%edi, %edi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%edx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	cudaSetDevice@PLT
	testl	%eax, %eax
	je	.L85
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$1, (%rax)
	movq	%rax, %rdi
	jmp	.L107
.L85:
	movl	%ebp, %ebx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	_Znam@PLT
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	movq	%rax, %r12
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L108
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L89
.L108:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$2, (%rax)
	movq	%rax, %rdi
	jmp	.L107
.L89:
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r14, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L91
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$3, (%rax)
	movq	%rax, %rdi
	jmp	.L107
.L91:
	movl	%ebp, %eax
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	shrl	$10, %eax
	salq	$10, %rdx
	movl	$1, %ecx
	incl	%eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L93
	movq	(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	%ebp, %ecx
	movl	%r13d, %edx
	call	_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij
.L93:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L94
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$4, (%rax)
	movq	%rax, %rdi
	jmp	.L107
.L94:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L96
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$5, (%rax)
	movq	%rax, %rdi
	jmp	.L107
.L96:
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	jne	.L106
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	je	.L84
.L106:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$6, (%rax)
	movq	%rax, %rdi
.L107:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L105
	xorl	%edx, %edx
	leaq	_ZTIi(%rip), %rsi
	call	__cxa_throw@PLT
.L84:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L102
.L105:
	call	__stack_chk_fail@PLT
.L102:
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2074:
	.size	_Z14gpuIntMultiplyPiij, .-_Z14gpuIntMultiplyPiij
	.globl	_Z12gpuIntDividePiij
	.type	_Z12gpuIntDividePiij, @function
_Z12gpuIntDividePiij:
.LFB2075:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	testl	%esi, %esi
	jne	.L110
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$7, (%rax)
	movq	%rax, %rdi
	jmp	.L134
.L110:
	movq	%rdi, %r14
	xorl	%edi, %edi
	movl	%esi, %r13d
	movl	%edx, %ebp
	call	cudaSetDevice@PLT
	testl	%eax, %eax
	je	.L112
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$1, (%rax)
	movq	%rax, %rdi
	jmp	.L134
.L112:
	movl	%ebp, %ebx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	_Znam@PLT
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	movq	%rax, %r12
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L135
	leaq	8(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L116
.L135:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$2, (%rax)
	movq	%rax, %rdi
	jmp	.L134
.L116:
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rdx
	movq	%r14, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L118
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$3, (%rax)
	movq	%rax, %rdi
	jmp	.L134
.L118:
	movl	%ebp, %eax
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	shrl	$10, %eax
	salq	$10, %rdx
	movl	$1, %ecx
	incl	%eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L120
	movq	(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	%ebp, %ecx
	movl	%r13d, %edx
	call	_Z36__device_stub__Z11intMultiplyPiPKiijPiPKiij
.L120:
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L121
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$4, (%rax)
	movq	%rax, %rdi
	jmp	.L134
.L121:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%r12, %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L123
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$5, (%rax)
	movq	%rax, %rdi
	jmp	.L134
.L123:
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	jne	.L133
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	testl	%eax, %eax
	je	.L109
.L133:
	movl	$4, %edi
	call	__cxa_allocate_exception@PLT
	movl	$6, (%rax)
	movq	%rax, %rdi
.L134:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L132
	xorl	%edx, %edx
	leaq	_ZTIi(%rip), %rsi
	call	__cxa_throw@PLT
.L109:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L129
.L132:
	call	__stack_chk_fail@PLT
.L129:
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2075:
	.size	_Z12gpuIntDividePiij, .-_Z12gpuIntDividePiij
	.globl	_Z33__device_stub__Z9intDividePiPKiijPiPKiij
	.type	_Z33__device_stub__Z9intDividePiPKiijPiPKiij, @function
_Z33__device_stub__Z9intDividePiPKiijPiPKiij:
.LFB2106:
	.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	.L136
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z9intDividePiPKiij(%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
.L136:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L138
	call	__stack_chk_fail@PLT
.L138:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2106:
	.size	_Z33__device_stub__Z9intDividePiPKiijPiPKiij, .-_Z33__device_stub__Z9intDividePiPKiijPiPKiij
	.globl	_Z9intDividePiPKiij
	.type	_Z9intDividePiPKiij, @function
_Z9intDividePiPKiij:
.LFB2107:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z9intDividePiPKiijPiPKiij
	.cfi_endproc
.LFE2107:
	.size	_Z9intDividePiPKiij, .-_Z9intDividePiPKiij
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z9intDividePiPKiij"
.LC1:
	.string	"_Z11intMultiplyPiPKiij"
.LC2:
	.string	"_Z11intSubtractPiPKiS1_j"
.LC3:
	.string	"_Z6intAddPiPKiS1_j"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2109:
	.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	_Z9intDividePiPKiij(%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	_Z11intMultiplyPiPKiij(%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	_Z11intSubtractPiPKiS1_j(%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	_Z6intAddPiPKiS1_j(%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
.LFE2109:
	.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:

#pragma once

#include "hip/hip_runtime.h"



#define MIN_GPU_ADDITION_ELEMENTS 500
#define MAX_ADDING_THREADS 1024

// Visable User Functions ---------------------

int* gpuIntAdd(int*, int*, unsigned int);
int* gpuIntSubtract(int*, int*, unsigned int);
int* gpuIntMultiply(int*, int, unsigned int);
int* gpuIntDivide(int*, int, unsigned int);


// GPU Functions ------------------------------

__global__ void intAdd(int*, const int*, const int*, const unsigned int);
__global__ void intSubtract(int*, const int*, const int*, const unsigned int);
__global__ void intMultiply(int*, const int*, const int, const unsigned int);
__global__ void intDivide(int*, const int*, const int, const unsigned int);
__global__ void floatAdd(float*, float*, float*);
__global__ void floatSubtract(float*, float*, float*);
__global__ void floatMultiply(float*, float*, float*);
__global__ void floatDivide(float*, float*, float*);
__global__ void doubleAdd(double*, double*, double*);
__global__ void doubleSubtract(double*, double*, double*);
__global__ void doubleMultiply(double*, double*, double*);
__global__ void doubleDivide(double*, double*, double*);
	.text
	.file	"LAVectorGPUOperations.hip"
	.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 __hip_cuid_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
