#include "GpuProjectionProcessing.cuh"
#include "GpuUtmWgsTransform.cuh"
#include "GpuTimer.cuh"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

using namespace winGpu;

__global__ void applyTransformUtmToWgsCoordsGpu(double xOrigin, double yOrigin, double xPixelSize, double yPixelSize,
	int height, int width, int zone, bool southhemi, double* lon, double* lat)
{
	int h = blockDim.x * blockIdx.x + threadIdx.x;
	int w = blockDim.y * blockIdx.y + threadIdx.y;
	if (height <= h || width <= w)
	{
		return;
	}
	double x = xOrigin + xPixelSize * w;
	double y = yOrigin + yPixelSize * h;
	double newLon = 0.0;
	double newLat = 0.0;

	UtmXYToLatLonGpu(x, y, zone, southhemi, newLon, newLat);

	lon[h * width + w] = newLon;
	lat[h * width + w] = newLat;
}

double winGpu::doTransformUtmToWgsCoordsGpu(double xOrigin, double yOrigin, double xPixelSize, double yPixelSize,
	int height, int width, int zone, bool southhemi, double* lon, double* lat)
{
	const size_t maxAvaliableCoords = 1500000;
	int countRowsPerIter = maxAvaliableCoords / width;
	int countIter = height / countRowsPerIter + 1;
	const size_t size = width * countRowsPerIter;
	dim3 threadsPerBlock(16, 16);
	dim3 numBlocks(countRowsPerIter / threadsPerBlock.x + 1, width / threadsPerBlock.y + 1);

	double* newLon = new double[size];
	double* newLat = new double[size];
	double* dev_lon = 0;
	double* dev_lat = 0;

	float time;
	cudaSetDevice(0);
	cudaMalloc((void**)&dev_lon, size * sizeof(double));
	cudaMalloc((void**)&dev_lat, size * sizeof(double));
	GPU_TIMER_START;
	for (int i = 0; i < countIter; i++)
	{
		double newYOrigin = yOrigin + i * yPixelSize * countRowsPerIter;

		applyTransformUtmToWgsCoordsGpu << <numBlocks, threadsPerBlock >> > (xOrigin, newYOrigin,
			xPixelSize, yPixelSize, countRowsPerIter, width, zone, southhemi, dev_lon, dev_lat);
		cudaDeviceSynchronize();

		cudaMemcpy(newLon, dev_lon, size * sizeof(double), cudaMemcpyDeviceToHost);
		cudaMemcpy(newLat, dev_lat, size * sizeof(double), cudaMemcpyDeviceToHost);

		size_t countCoordsForCopy = i != countIter - 1 ? size :
			width * height - countRowsPerIter * width * i;
		for (int j = 0; j < countCoordsForCopy; j++)
		{
			lon[i * size + j] = newLon[j];
			lat[i * size + j] = newLat[j];
		}
	}
	GPU_TIMER_STOP(time);
	cudaFree(dev_lon);
	cudaFree(dev_lat);
	delete[] newLon;
	delete[] newLat;

	return (double)time;
}

__global__ void applyTransformWgsToUtmCoordsGpu(double xOrigin, double yOrigin, double xPixelSize, double yPixelSize,
	int height, int width, int zone, double* x, double* y)
{
	int h = blockDim.x * blockIdx.x + threadIdx.x;
	int w = blockDim.y * blockIdx.y + threadIdx.y;
	if (height <= h || width <= w)
	{
		return;
	}
	double lon = xOrigin + xPixelSize * w;
	double lat = yOrigin + yPixelSize * h;
	double newX = 0.0;
	double newY = 0.0;

	LatLonToUtmXYGpu(lon, lat, zone, newX, newY);

	x[h * width + w] = newX;
	y[h * width + w] = newY;
}

double winGpu::doTransformWgsToUtmCoordsGpu(double xOrigin, double yOrigin, double xPixelSize, double yPixelSize,
	int height, int width, int zone, double* x, double* y)
{
	const size_t maxAvaliableCoords = 2000000;
	int countRowsPerIter = maxAvaliableCoords / width;
	int countIter = height / countRowsPerIter + 1;
	const size_t size = width * countRowsPerIter;
	dim3 threadsPerBlock(16, 16);
	dim3 numBlocks(countRowsPerIter / threadsPerBlock.x + 1, width / threadsPerBlock.y + 1);

	double* newX = new double[size];
	double* newY = new double[size];
	double* dev_x = 0;
	double* dev_y = 0;

	float time;
	cudaSetDevice(0);
	cudaMalloc((void**)&dev_x, size * sizeof(double));
	cudaMalloc((void**)&dev_y, size * sizeof(double));
	GPU_TIMER_START;
	for (int i = 0; i < countIter; i++)
	{
		double newYOrigin = yOrigin + i * yPixelSize * countRowsPerIter;

		applyTransformWgsToUtmCoordsGpu << <numBlocks, threadsPerBlock >> > (xOrigin, newYOrigin,
			xPixelSize, yPixelSize, countRowsPerIter, width, zone, dev_x, dev_y);
		cudaDeviceSynchronize();

		cudaMemcpy(newX, dev_x, size * sizeof(double), cudaMemcpyDeviceToHost);
		cudaMemcpy(newY, dev_y, size * sizeof(double), cudaMemcpyDeviceToHost);

		size_t countCoordsForCopy = i != countIter - 1 ? size :
			width * height - countRowsPerIter * width * i;
		for (int j = 0; j < countCoordsForCopy; j++)
		{
			x[i * size + j] = newX[j];
			y[i * size + j] = newY[j];
		}
	}
	GPU_TIMER_STOP(time);
	cudaFree(dev_x);
	cudaFree(dev_y);
	delete[] newX;
	delete[] newY;

	return (double)time;
}

	.file	"tmpxft_002e2744_00000000-6_GpuProjectionProcessing.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2050:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2050:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0, @function
_ZN7mathGpu8floorGpuEd.constprop.0.isra.0:
.LFB2216:
	.cfi_startproc
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2216:
	.size	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0, .-_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.globl	_ZN7mathGpu8floorGpuEd
	.type	_ZN7mathGpu8floorGpuEd, @function
_ZN7mathGpu8floorGpuEd:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_ZN7mathGpu8floorGpuEd, .-_ZN7mathGpu8floorGpuEd
	.globl	_ZN7mathGpu6absGpuEd
	.type	_ZN7mathGpu6absGpuEd, @function
_ZN7mathGpu6absGpuEd:
.LFB2194:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2194:
	.size	_ZN7mathGpu6absGpuEd, .-_ZN7mathGpu6absGpuEd
	.globl	_ZN7mathGpu9newtonGpuEdd
	.type	_ZN7mathGpu9newtonGpuEdd, @function
_ZN7mathGpu9newtonGpuEdd:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2029:
	.size	_ZN7mathGpu9newtonGpuEdd, .-_ZN7mathGpu9newtonGpuEdd
	.globl	_ZN7mathGpu7sqrtGpuEd
	.type	_ZN7mathGpu7sqrtGpuEd, @function
_ZN7mathGpu7sqrtGpuEd:
.LFB2196:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2196:
	.size	_ZN7mathGpu7sqrtGpuEd, .-_ZN7mathGpu7sqrtGpuEd
	.globl	_ZN7mathGpu7factGpuEi
	.type	_ZN7mathGpu7factGpuEi, @function
_ZN7mathGpu7factGpuEi:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZN7mathGpu7factGpuEi, .-_ZN7mathGpu7factGpuEi
	.globl	_ZN7mathGpu6powGpuEdi
	.type	_ZN7mathGpu6powGpuEdi, @function
_ZN7mathGpu6powGpuEdi:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZN7mathGpu6powGpuEdi, .-_ZN7mathGpu6powGpuEdi
	.globl	_ZN7mathGpu14simplifyRadGpuEd
	.type	_ZN7mathGpu14simplifyRadGpuEd, @function
_ZN7mathGpu14simplifyRadGpuEd:
.LFB2198:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2198:
	.size	_ZN7mathGpu14simplifyRadGpuEd, .-_ZN7mathGpu14simplifyRadGpuEd
	.globl	_ZN7mathGpu6cosGpuEdi
	.type	_ZN7mathGpu6cosGpuEdi, @function
_ZN7mathGpu6cosGpuEdi:
.LFB2208:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2208:
	.size	_ZN7mathGpu6cosGpuEdi, .-_ZN7mathGpu6cosGpuEdi
	.globl	_ZN7mathGpu6sinGpuEdi
	.type	_ZN7mathGpu6sinGpuEdi, @function
_ZN7mathGpu6sinGpuEdi:
.LFB2210:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2210:
	.size	_ZN7mathGpu6sinGpuEdi, .-_ZN7mathGpu6sinGpuEdi
	.globl	_ZN7mathGpu6tanGpuEdi
	.type	_ZN7mathGpu6tanGpuEdi, @function
_ZN7mathGpu6tanGpuEdi:
.LFB2212:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2212:
	.size	_ZN7mathGpu6tanGpuEdi, .-_ZN7mathGpu6tanGpuEdi
	.globl	_ZN6winGpu11DegToRadGpuEd
	.type	_ZN6winGpu11DegToRadGpuEd, @function
_ZN6winGpu11DegToRadGpuEd:
.LFB2200:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2200:
	.size	_ZN6winGpu11DegToRadGpuEd, .-_ZN6winGpu11DegToRadGpuEd
	.globl	_ZN6winGpu11RadToDegGpuEd
	.type	_ZN6winGpu11RadToDegGpuEd, @function
_ZN6winGpu11RadToDegGpuEd:
.LFB2202:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2202:
	.size	_ZN6winGpu11RadToDegGpuEd, .-_ZN6winGpu11RadToDegGpuEd
	.globl	_ZN6winGpu22ArcLengthOfMeridianGpuEd
	.type	_ZN6winGpu22ArcLengthOfMeridianGpuEd, @function
_ZN6winGpu22ArcLengthOfMeridianGpuEd:
.LFB2204:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2204:
	.size	_ZN6winGpu22ArcLengthOfMeridianGpuEd, .-_ZN6winGpu22ArcLengthOfMeridianGpuEd
	.globl	_ZN6winGpu21UTMCentralMeridianGpuEi
	.type	_ZN6winGpu21UTMCentralMeridianGpuEi, @function
_ZN6winGpu21UTMCentralMeridianGpuEi:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2040:
	.size	_ZN6winGpu21UTMCentralMeridianGpuEi, .-_ZN6winGpu21UTMCentralMeridianGpuEi
	.globl	_ZN6winGpu20FootpointLatitudeGpuEd
	.type	_ZN6winGpu20FootpointLatitudeGpuEd, @function
_ZN6winGpu20FootpointLatitudeGpuEd:
.LFB2206:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN7mathGpu8floorGpuEd.constprop.0.isra.0
	.cfi_endproc
.LFE2206:
	.size	_ZN6winGpu20FootpointLatitudeGpuEd, .-_ZN6winGpu20FootpointLatitudeGpuEd
	.globl	_ZN6winGpu16MapLatLonToXYGpuEdddRdS0_
	.type	_ZN6winGpu16MapLatLonToXYGpuEdddRdS0_, @function
_ZN6winGpu16MapLatLonToXYGpuEdddRdS0_:
.LFB2042:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2042:
	.size	_ZN6winGpu16MapLatLonToXYGpuEdddRdS0_, .-_ZN6winGpu16MapLatLonToXYGpuEdddRdS0_
	.globl	_ZN6winGpu16MapXYToLatLonGpuEdddRdS0_
	.type	_ZN6winGpu16MapXYToLatLonGpuEdddRdS0_, @function
_ZN6winGpu16MapXYToLatLonGpuEdddRdS0_:
.LFB2214:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2214:
	.size	_ZN6winGpu16MapXYToLatLonGpuEdddRdS0_, .-_ZN6winGpu16MapXYToLatLonGpuEdddRdS0_
	.globl	_ZN6winGpu16LatLonToUtmXYGpuEddiRdS0_
	.type	_ZN6winGpu16LatLonToUtmXYGpuEddiRdS0_, @function
_ZN6winGpu16LatLonToUtmXYGpuEddiRdS0_:
.LFB2044:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2044:
	.size	_ZN6winGpu16LatLonToUtmXYGpuEddiRdS0_, .-_ZN6winGpu16LatLonToUtmXYGpuEddiRdS0_
	.globl	_ZN6winGpu16UtmXYToLatLonGpuEddibRdS0_
	.type	_ZN6winGpu16UtmXYToLatLonGpuEddibRdS0_, @function
_ZN6winGpu16UtmXYToLatLonGpuEddibRdS0_:
.LFB2045:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2045:
	.size	_ZN6winGpu16UtmXYToLatLonGpuEddibRdS0_, .-_ZN6winGpu16UtmXYToLatLonGpuEddibRdS0_
	.globl	_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_
	.type	_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_, @function
_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movl	%edi, 28(%rsp)
	leaq	88(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	100(%rsp), %rsi
	movl	%edx, 20(%rsp)
	leaq	72(%rsp), %rdx
	movb	%cl, 16(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%r9, (%rsp)
	movsd	%xmm0, 56(%rsp)
	movsd	%xmm1, 48(%rsp)
	movsd	%xmm2, 40(%rsp)
	movsd	%xmm3, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 200(%rsp)
	movq	%rsp, %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L42
	pushq	80(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 256
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L42:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L44
	call	__stack_chk_fail@PLT
.L44:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_, .-_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_
	.globl	_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_
	.type	_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_, @function
_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_:
.LFB2073:
	.cfi_startproc
	endbr64
	movzbl	%cl, %ecx
	jmp	_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_
	.cfi_endproc
.LFE2073:
	.size	_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_, .-_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_
	.globl	_ZN6winGpu28doTransformUtmToWgsCoordsGpuEddddiiibPdS0_
	.type	_ZN6winGpu28doTransformUtmToWgsCoordsGpuEddddiiibPdS0_, @function
_ZN6winGpu28doTransformUtmToWgsCoordsGpuEddddiiibPdS0_:
.LFB2046:
	.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	$184, %rsp
	.cfi_def_cfa_offset 240
	movl	%edx, 72(%rsp)
	xorl	%edx, %edx
	movb	%cl, 76(%rsp)
	movslq	%esi, %rcx
	movq	%r8, 80(%rsp)
	movq	%r9, 88(%rsp)
	movsd	%xmm0, 48(%rsp)
	movsd	%xmm1, 56(%rsp)
	movsd	%xmm2, 64(%rsp)
	movsd	%xmm3, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	movl	$1500000, %eax
	movl	$1, 152(%rsp)
	divq	%rcx
	movl	$1, 164(%rsp)
	movq	%rax, %rcx
	movl	%eax, 20(%rsp)
	movl	%edi, %eax
	cltd
	idivl	%ecx
	movl	%eax, 44(%rsp)
	movl	%esi, %eax
	imull	%ecx, %eax
	shrl	$4, %ecx
	incl	%ecx
	movl	%ecx, 156(%rsp)
	movl	%eax, 40(%rsp)
	movslq	%eax, %r13
	movl	%esi, %eax
	shrl	$4, %eax
	incl	%eax
	movl	%eax, 160(%rsp)
	movabsq	$1152921504606846975, %rax
	cmpq	%r13, %rax
	jb	.L49
	leaq	0(,%r13,8), %rax
	movl	%edi, %ebx
	movl	%esi, %ebp
	xorl	%r14d, %r14d
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	imull	%ebp, %ebx
	xorl	%r12d, %r12d
	call	_Znam@PLT
	leaq	0(,%r13,8), %rdi
	movq	%rax, 24(%rsp)
	call	_Znam@PLT
	xorl	%edi, %edi
	movq	%rax, %r15
	xorl	%eax, %eax
	movq	%rax, 112(%rsp)
	movq	%rax, 120(%rsp)
	call	cudaSetDevice@PLT
	movq	8(%rsp), %rsi
	leaq	112(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	leaq	120(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	128(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	136(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movzbl	76(%rsp), %eax
	movl	%eax, 76(%rsp)
	jmp	.L50
.L49:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L60
	call	__cxa_throw_bad_array_new_length@PLT
.L62:
	movl	40(%rsp), %eax
	incl	%r12d
	addq	%r13, %r14
	subl	%eax, %ebx
.L50:
	cmpl	%r12d, 44(%rsp)
	jl	.L61
	movl	152(%rsp), %ecx
	movl	164(%rsp), %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	156(%rsp), %rdi
	movabsq	$68719476752, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L52
	cvtsi2sdl	%r12d, %xmm1
	movq	120(%rsp), %r9
	movl	76(%rsp), %ecx
	movl	%ebp, %esi
	mulsd	32(%rsp), %xmm1
	cvtsi2sdl	20(%rsp), %xmm0
	movq	112(%rsp), %r8
	movl	72(%rsp), %edx
	movl	20(%rsp), %edi
	movsd	32(%rsp), %xmm3
	movsd	64(%rsp), %xmm2
	mulsd	%xmm0, %xmm1
	movsd	48(%rsp), %xmm0
	addsd	56(%rsp), %xmm1
	call	_Z61__device_stub__Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_ddddiiibPdS_
.L52:
	call	cudaDeviceSynchronize@PLT
	movq	8(%rsp), %rdx
	movq	112(%rsp), %rsi
	movl	$2, %ecx
	movq	24(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rsi
	movq	%r15, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movq	%r13, %rcx
	cmpl	%r12d, 44(%rsp)
	jne	.L53
	movslq	%ebx, %rcx
.L53:
	movq	80(%rsp), %rdx
	leaq	0(,%r14,8), %rax
	leaq	(%rdx,%rax), %rsi
	movq	88(%rsp), %rdx
	addq	%rax, %rdx
	xorl	%eax, %eax
.L54:
	cmpq	%rax, %rcx
	je	.L62
	movq	24(%rsp), %rdi
	movsd	(%rdi,%rax,8), %xmm0
	movsd	%xmm0, (%rsi,%rax,8)
	movsd	(%r15,%rax,8), %xmm0
	movsd	%xmm0, (%rdx,%rax,8)
	incq	%rax
	jmp	.L54
.L61:
	movq	136(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	136(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	136(%rsp), %rdx
	leaq	108(%rsp), %rdi
	movq	128(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	_ZdaPv@PLT
	movq	%r15, %rdi
	call	_ZdaPv@PLT
	cvtss2sd	108(%rsp), %xmm0
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L57
.L60:
	call	__stack_chk_fail@PLT
.L57:
	addq	$184, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2046:
	.size	_ZN6winGpu28doTransformUtmToWgsCoordsGpuEddddiiibPdS0_, .-_ZN6winGpu28doTransformUtmToWgsCoordsGpuEddddiiibPdS0_
	.globl	_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_
	.type	_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_, @function
_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_:
.LFB2074:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movl	%edi, 28(%rsp)
	leaq	80(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	92(%rsp), %rsi
	movl	%edx, 20(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, (%rsp)
	movsd	%xmm0, 56(%rsp)
	movsd	%xmm1, 48(%rsp)
	movsd	%xmm2, 40(%rsp)
	movsd	%xmm3, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L63
	pushq	72(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 240
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L63:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L65
	call	__stack_chk_fail@PLT
.L65:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2074:
	.size	_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_, .-_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_
	.globl	_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_
	.type	_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_, @function
_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_:
.LFB2075:
	.cfi_startproc
	endbr64
	jmp	_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_
	.cfi_endproc
.LFE2075:
	.size	_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_, .-_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_
	.globl	_ZN6winGpu28doTransformWgsToUtmCoordsGpuEddddiiiPdS0_
	.type	_ZN6winGpu28doTransformWgsToUtmCoordsGpuEddddiiiPdS0_, @function
_ZN6winGpu28doTransformWgsToUtmCoordsGpuEddddiiiPdS0_:
.LFB2047:
	.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	$168, %rsp
	.cfi_def_cfa_offset 224
	movl	%edx, 60(%rsp)
	xorl	%edx, %edx
	movq	%rcx, 64(%rsp)
	movslq	%esi, %rcx
	movq	%r8, 72(%rsp)
	movsd	%xmm0, 32(%rsp)
	movsd	%xmm1, 40(%rsp)
	movsd	%xmm2, 48(%rsp)
	movsd	%xmm3, 24(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	movl	$2000000, %eax
	movl	$1, 136(%rsp)
	divq	%rcx
	movl	$1, 148(%rsp)
	movq	%rax, %rcx
	movl	%eax, 12(%rsp)
	movl	%edi, %eax
	cltd
	idivl	%ecx
	movl	%eax, 56(%rsp)
	movl	%esi, %eax
	imull	%ecx, %eax
	shrl	$4, %ecx
	incl	%ecx
	movl	%ecx, 140(%rsp)
	movl	%eax, 16(%rsp)
	movslq	%eax, %r14
	movl	%esi, %eax
	shrl	$4, %eax
	incl	%eax
	movl	%eax, 144(%rsp)
	movabsq	$1152921504606846975, %rax
	cmpq	%r14, %rax
	jb	.L69
	leaq	0(,%r14,8), %rax
	movl	%edi, %r13d
	movl	%esi, %ebx
	xorl	%r15d, %r15d
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	imull	%ebx, %r13d
	call	_Znam@PLT
	leaq	0(,%r14,8), %rdi
	movq	%rax, %r12
	call	_Znam@PLT
	xorl	%edi, %edi
	movq	%rax, %rbp
	xorl	%eax, %eax
	movq	%rax, 96(%rsp)
	movq	%rax, 104(%rsp)
	call	cudaSetDevice@PLT
	movq	(%rsp), %rsi
	leaq	96(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	104(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	112(%rsp), %rdi
	call	cudaEventCreate@PLT
	leaq	120(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	%r13d, 20(%rsp)
	xorl	%r13d, %r13d
	jmp	.L70
.L69:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L80
	call	__cxa_throw_bad_array_new_length@PLT
.L82:
	movl	16(%rsp), %edx
	subl	%edx, 20(%rsp)
	incl	%r13d
	addq	%r14, %r15
.L70:
	cmpl	%r13d, 56(%rsp)
	jl	.L81
	movl	136(%rsp), %ecx
	movl	148(%rsp), %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	140(%rsp), %rdi
	movabsq	$68719476752, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L72
	cvtsi2sdl	%r13d, %xmm1
	movq	104(%rsp), %r8
	movl	60(%rsp), %edx
	movl	%ebx, %esi
	mulsd	24(%rsp), %xmm1
	cvtsi2sdl	12(%rsp), %xmm0
	movq	96(%rsp), %rcx
	movl	12(%rsp), %edi
	movsd	24(%rsp), %xmm3
	movsd	48(%rsp), %xmm2
	mulsd	%xmm0, %xmm1
	movsd	32(%rsp), %xmm0
	addsd	40(%rsp), %xmm1
	call	_Z60__device_stub__Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_ddddiiiPdS_
.L72:
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	%r12, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	%rbp, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movq	%r14, %rcx
	cmpl	%r13d, 56(%rsp)
	jne	.L73
	movslq	20(%rsp), %rcx
.L73:
	movq	64(%rsp), %rdx
	leaq	0(,%r15,8), %rax
	leaq	(%rdx,%rax), %rsi
	movq	72(%rsp), %rdx
	addq	%rax, %rdx
	xorl	%eax, %eax
.L74:
	cmpq	%rax, %rcx
	je	.L82
	movsd	(%r12,%rax,8), %xmm0
	movsd	%xmm0, (%rsi,%rax,8)
	movsd	0(%rbp,%rax,8), %xmm0
	movsd	%xmm0, (%rdx,%rax,8)
	incq	%rax
	jmp	.L74
.L81:
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	120(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	leaq	92(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	movq	%rbp, %rdi
	call	_ZdaPv@PLT
	cvtss2sd	92(%rsp), %xmm0
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L77
.L80:
	call	__stack_chk_fail@PLT
.L77:
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2047:
	.size	_ZN6winGpu28doTransformWgsToUtmCoordsGpuEddddiiiPdS0_, .-_ZN6winGpu28doTransformWgsToUtmCoordsGpuEddddiiiPdS0_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_"
.LC1:
	.string	"_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2077:
	.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	_Z31applyTransformWgsToUtmCoordsGpuddddiiiPdS_(%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	_Z31applyTransformUtmToWgsCoordsGpuddddiiibPdS_(%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
.LFE2077:
	.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"

#include <stdio.h>

typedef short int pixel;

namespace winGpu
{
	double doTransformUtmToWgsCoordsGpu(double xOrigin, double yOrigin, double xPixelSize, double yPixelSize,
		int height, int width, int zone, bool southhemi, double* lon, double* lat);
	double doTransformWgsToUtmCoordsGpu(double xOrigin, double yOrigin, double xPixelSize, double yPixelSize,
		int height, int width, int zone, double* x, double* y);
}
	.text
	.file	"GpuProjectionProcessing.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
