﻿#include "ZonePlanMCMC.cuh"
#include <vector>
#include <iostream>

__device__
__host__
unsigned int rand(unsigned int* randx) {
    *randx = *randx * 1103515245 + 12345;
    return (*randx)&2147483647;
}

__device__
__host__
float randf(unsigned int* randx) {
	return rand(randx) / (float(2147483647) + 1);
}

__device__
__host__
float randf(unsigned int* randx, float a, float b) {
	return randf(randx) * (b - a) + a;
}

__device__
__host__
float sampleFromCdf(unsigned int* randx, float* cdf, int num) {
	float rnd = randf(randx, 0, cdf[num-1]);

	for (int i = 0; i < num; ++i) {
		if (rnd <= cdf[i]) return i;
	}

	return num - 1;
}

__device__
__host__
float sampleFromPdf(unsigned int* randx, float* pdf, int num) {
	if (num == 0) return 0;

	float cdf[40];
	cdf[0] = pdf[0];
	for (int i = 1; i < num; ++i) {
		if (cdf[i] >= 0) {
			cdf[i] = cdf[i - 1] + pdf[i];
		} else {
			cdf[i] = cdf[i - 1];
		}
	}

	return sampleFromCdf(randx, cdf, num);
}

__device__
__host__
float sum(float* values, int num) {
	float total = 0.0f;
	for (int i = 0; i < num; ++i) {
		total += values[i];
	}
	return total;
}

__device__
__host__
void swapZoneType(zone_type* z1, zone_type* z2) {
	int temp_type = z1->type;
	int temp_level = z1->level;

	z1->type = z2->type;
	z1->level = z2->level;
	z2->type = temp_type;
	z2->level = temp_level;
}

__device__
__host__
float dot2(float* preference, float* feature) {
	float ret = 0.0;
	for (int i = 0; i < 9; ++i) {
		ret += preference[i] * feature[i];
	}
	return ret;
}

__device__
__host__
float noise(float distToFactory, float distToAmusement, float distToStore) {
	float Km = 800.0 - distToFactory;
	float Ka = 400.0 - distToAmusement;
	float Ks = 200.0 - distToStore;

	return max(max(max(Km, Ka), Ks), 0.0f);
}

__device__
__host__
float pollution(float distToFactory) {
	float Km = 800.0 - distToFactory;

	return max(Km, 0.0f);
}

__device__
__host__
void MCMCstep(int numIterations, unsigned int randx, zone_plan* plan, zone_plan* proposal, zone_plan* bestPlan) {
	float K = 0.001;

	// preference
	float preference[10][9];
	preference[0][0] = 0; preference[0][1] = 0; preference[0][2] = 0.15; preference[0][3] = 0.15; preference[0][4] = 0.3; preference[0][5] = 0; preference[0][6] = 0.1; preference[0][7] = 0.1; preference[0][8] = 0.2;
	preference[1][0] = 0; preference[1][1] = 0; preference[1][2] = 0.15; preference[1][3] = 0; preference[1][4] = 0.55; preference[1][5] = 0; preference[1][6] = 0.2; preference[1][7] = 0.1; preference[1][8] = 0;
	preference[2][0] = 0; preference[2][1] = 0; preference[2][2] = 0.05; preference[2][3] = 0; preference[2][4] = 0; preference[2][5] = 0; preference[2][6] = 0.25; preference[2][7] = 0.1; preference[2][8] = 0.6;
	preference[3][0] = 0.18; preference[3][1] = 0.17; preference[3][2] = 0; preference[3][3] = 0.17; preference[3][4] = 0; preference[3][5] = 0.08; preference[3][6] = 0.2; preference[3][7] = 0.2; preference[3][8] = 0;
	preference[4][0] = 0.3; preference[4][1] = 0; preference[4][2] = 0.3; preference[4][3] = 0.1; preference[4][4] = 0; preference[4][5] = 0; preference[4][6] = 0.1; preference[4][7] = 0.2; preference[4][8] = 0;
	preference[5][0] = 0.05; preference[5][1] = 0; preference[5][2] = 0.1; preference[5][3] = 0.2; preference[5][4] = 0.1; preference[5][5] = 0; preference[5][6] = 0.1; preference[5][7] = 0.15; preference[5][8] = 0.3;
	preference[6][0] = 0.15; preference[6][1] = 0.1; preference[6][2] = 0; preference[6][3] = 0.15; preference[6][4] = 0; preference[6][5] = 0.1; preference[6][6] = 0.1; preference[6][7] = 0.2; preference[6][8] = 0.2;
	preference[7][0] = 0.2; preference[7][1] = 0; preference[7][2] = 0.25; preference[7][3] = 0; preference[7][4] = 0.15; preference[7][5] = 0; preference[7][6] = 0.1; preference[7][7] = 0.1; preference[7][8] = 0.2;
	preference[8][0] = 0.3; preference[8][1] = 0; preference[8][2] = 0.15; preference[8][3] = 0.05; preference[8][4] = 0; preference[8][5] = 0; preference[8][6] = 0.25; preference[8][7] = 0.25; preference[8][8] = 0;
	preference[9][0] = 0.4; preference[9][1] = 0; preference[9][2] = 0.2; preference[9][3] = 0; preference[9][4] = 0; preference[9][5] = 0; preference[9][6] = 0.2; preference[9][7] = 0.2; preference[9][8] = 0;

	// population ratio
	float ratioPeople[10] = {0.06667, 0.06667, 0.06667, 0.21, 0.09, 0.09, 0.09, 0.12, 0.1, 0.1};

	// poppulation of each level of zone
	float levelPeople[3] = {1, 5, 10};

	// initial plan
	{
		float zoneTypeDistribution[18] = {0.2, 0.38, 0.2, 0.06, 0.05, 0.03, 0.02, 0.01, 0.01, 0.02, 0, 0, 0.01, 0, 0, 0.01, 0, 0};
		float Z = sum(zoneTypeDistribution, 18);

		float remainedBlockNum[18];
		for (int zi = 0; zi < 18; ++zi) {
			remainedBlockNum[zi] = zoneTypeDistribution[zi] / Z * ZONE_GRID_SIZE * ZONE_GRID_SIZE;
		}

		for (int r = 0; r < ZONE_GRID_SIZE; ++r) {
			for (int c = 0; c < ZONE_GRID_SIZE; ++c) {
				int n = sampleFromPdf(&randx, remainedBlockNum, 18);
				plan->zones[r][c].type = n / 3;
				plan->zones[r][c].level = n % 3 + 1;
				remainedBlockNum[n]--;
			}
		}

		plan->score = 0.0;
	}

	bestPlan->score = 0.0;
	//float current_score = 0.0f;
	for (int loop = 0; loop < numIterations; ++loop) {
		// create a proposal
		{
			// copy the current plan to the proposal
			*proposal = *plan;

			// swap a zone type between two blocks
			while (true) {
				int x1 = randf(&randx, 0, ZONE_GRID_SIZE);
				int y1 = randf(&randx, 0, ZONE_GRID_SIZE);
				int x2 = randf(&randx, 0, ZONE_GRID_SIZE);
				int y2 = randf(&randx, 0, ZONE_GRID_SIZE);

				if (proposal->zones[y1][x1].type != proposal->zones[y2][x2].type || proposal->zones[y1][x1].level != proposal->zones[y2][x2].level) {
					swapZoneType(&proposal->zones[y1][x1], &proposal->zones[y2][x2]);
					break;
				}
			}
		}

		// 
		proposal->score = 0.0;
		float count = 0.0;

		for (int r = 0; r < ZONE_GRID_SIZE; ++r) {
			for (int c = 0; c < ZONE_GRID_SIZE; ++c) {
				// skip for non-residential block
				if (plan->zones[r][c].type != 0) continue;

				// compute the distance to the nearest spots
				float distToStore = 4000;
				float distToRestaurant = 4000;
				float distToFactory = 4000;
				float distToPark = 4000;
				float distToAmusement = 4000;
				float distToSchool = 4000;
				float distToLibrary = 4000;

				for (int r2 = 0; r2 < ZONE_GRID_SIZE; ++r2) {
					for (int c2 = 0; c2 < ZONE_GRID_SIZE; ++c2) {
						if (proposal->zones[r2][c2].type == 0) continue;

						//float dist = ZONE_CELL_LEN * sqrtf((r - r2) * (r - r2) + (c - c2) * (c - c2));
						float dist = ZONE_CELL_LEN * (abs(r - r2) + abs(c - c2));
						if (proposal->zones[r2][c2].type == 1) { // 店・レストラン
							if (dist < distToStore) {
								distToStore = dist;
								distToRestaurant = dist;
							}
						} else if (proposal->zones[r2][c2].type == 2) { // 工場
							if (dist < distToFactory) {
								distToFactory = dist;
							}
						} else if (proposal->zones[r2][c2].type == 3) { // 公園
							if (dist < distToPark) {
								distToPark = dist;
							}
						} else if (proposal->zones[r2][c2].type == 4) { // アミューズメント
							if (dist < distToAmusement) {
								distToAmusement = dist;
							}
						} else if (proposal->zones[r2][c2].type == 5) { // 学校・図書館
							if (dist < distToSchool) {
								distToSchool = dist;
								distToLibrary = dist;
							}
						}
					}
				}

				// compute feature
				float feature[9];
				feature[0] = expf(-K * distToStore);
				feature[1] = expf(-K * distToSchool);
				feature[2] = expf(-K * distToRestaurant);
				feature[3] = expf(-K * distToPark);
				feature[4] = expf(-K * distToAmusement);
				feature[5] = expf(-K * distToLibrary);
				feature[6] = expf(-K * noise(distToFactory, distToAmusement, distToStore));
				feature[7] = expf(-K * pollution(distToFactory));
				feature[8] = 0; // 駅はなし

				// compute score
				for (int i = 0; i < 10; ++i) {
					proposal->score += dot2(preference[i], feature) * ratioPeople[i] * levelPeople[proposal->zones[r][c].level - 1];
					count += ratioPeople[i] * levelPeople[proposal->zones[r][c].level - 1];
				}
			}
		}

		proposal->score /= count;

		// update the best plan
		if (proposal->score > bestPlan->score) {
			*bestPlan = *proposal;
		}

		// compare the current plan and the proposal
		if (proposal->score > plan->score || loop % 10 == 0) {
			// accept
			*plan = *proposal;
		}
	}
}

/**
 * CUDA version of MCMCM
 */
__global__
void zonePlanMCMCGPUKernel(int* numIterations, zone_plan* plan, zone_plan* proposal, zone_plan* bestPlan) {
	int idx = blockDim.x * blockIdx.x + threadIdx.x;

	// initialize random
	unsigned int randx = idx;

	MCMCstep(*numIterations, randx, &plan[idx], &proposal[idx], &bestPlan[idx]);
}

/**
 * CUDA version of MCMC
 */
void zonePlanMCMCGPUfunc(zone_plan** bestPlans, int numIterations) {
	// CPU側のメモリを確保
	*bestPlans = (zone_plan*)malloc(sizeof(zone_plan) * ZONE_PLAN_MCMC_GRID_SIZE * ZONE_PLAN_MCMC_BLOCK_SIZE);

	// デバイスメモリを確保
	int* devNumIterations;
	if (cudaMalloc((void**)&devNumIterations, sizeof(int)) != cudaSuccess) {
		printf("cuda memory allocation error!\n");
		return;
	}
	zone_plan* devPlan;
	if (cudaMalloc((void**)&devPlan, sizeof(zone_plan) * ZONE_PLAN_MCMC_GRID_SIZE * ZONE_PLAN_MCMC_BLOCK_SIZE) != cudaSuccess) {
		printf("cuda memory allocation error!\n");
		return;
	}
	zone_plan* devProposal;
	if (cudaMalloc((void**)&devProposal, sizeof(zone_plan) * ZONE_PLAN_MCMC_GRID_SIZE * ZONE_PLAN_MCMC_BLOCK_SIZE) != cudaSuccess) {
		cudaFree(devPlan);
		printf("cuda memory allocation error!\n");
		return;
	}
	zone_plan* devBestPlan;
	if (cudaMalloc((void**)&devBestPlan, sizeof(zone_plan) * ZONE_PLAN_MCMC_GRID_SIZE * ZONE_PLAN_MCMC_BLOCK_SIZE) != cudaSuccess) {
		cudaFree(devPlan);
		cudaFree(devProposal);
		printf("cuda memory allocation error!\n");
		return;
	}

	// copy memory
	if (cudaMemcpy(devNumIterations, &numIterations, sizeof(int), cudaMemcpyHostToDevice) != cudaSuccess) {
		cudaFree(devPlan);
		cudaFree(devProposal);
		cudaFree(devBestPlan);
		printf("cuda memory copy error!\n");
		return;
	}

	printf("start GPU kernel.\n");

	// GPU側のカーネル関数を呼び出す
    zonePlanMCMCGPUKernel<<<ZONE_PLAN_MCMC_GRID_SIZE, ZONE_PLAN_MCMC_BLOCK_SIZE>>>(devNumIterations, devPlan, devProposal, devBestPlan);

	// 結果をCPU側のバッファへ転送する
	if (cudaMemcpy(*bestPlans, devBestPlan, sizeof(zone_plan) * ZONE_PLAN_MCMC_GRID_SIZE * ZONE_PLAN_MCMC_BLOCK_SIZE, cudaMemcpyDeviceToHost) != cudaSuccess) {
		cudaFree(devPlan);
		cudaFree(devProposal);
		cudaFree(devBestPlan);
		printf("cuda memory copy error!\n");
		return;
	}

	// デバイスメモリを開放する
    cudaFree(devPlan);
    cudaFree(devProposal);
    cudaFree(devBestPlan);

	printf("GPU kernel done.\n");
}

	.file	"tmpxft_0033f361_00000000-6_ZonePlanMCMC.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4012:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4012:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z4randPj
	.type	_Z4randPj, @function
_Z4randPj:
.LFB3998:
	.cfi_startproc
	endbr64
	imull	$1103515245, (%rdi), %eax
	addl	$12345, %eax
	movl	%eax, (%rdi)
	andl	$2147483647, %eax
	ret
	.cfi_endproc
.LFE3998:
	.size	_Z4randPj, .-_Z4randPj
	.globl	_Z5randfPj
	.type	_Z5randfPj, @function
_Z5randfPj:
.LFB3999:
	.cfi_startproc
	endbr64
	call	_Z4randPj
	movl	%eax, %eax
	cvtsi2ssq	%rax, %xmm0
	mulss	.LC2(%rip), %xmm0
	ret
	.cfi_endproc
.LFE3999:
	.size	_Z5randfPj, .-_Z5randfPj
	.globl	_Z5randfPjff
	.type	_Z5randfPjff, @function
_Z5randfPjff:
.LFB4000:
	.cfi_startproc
	endbr64
	movaps	%xmm0, %xmm2
	call	_Z5randfPj
	subss	%xmm2, %xmm1
	mulss	%xmm0, %xmm1
	addss	%xmm2, %xmm1
	movaps	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE4000:
	.size	_Z5randfPjff, .-_Z5randfPjff
	.globl	_Z13sampleFromCdfPjPfi
	.type	_Z13sampleFromCdfPjPfi, @function
_Z13sampleFromCdfPjPfi:
.LFB4001:
	.cfi_startproc
	endbr64
	movslq	%edx, %rax
	xorps	%xmm0, %xmm0
	movss	-4(%rsi,%rax,4), %xmm1
	call	_Z5randfPjff
	xorl	%eax, %eax
.L8:
	cmpl	%eax, %edx
	jle	.L16
	leaq	1(%rax), %rcx
	movss	-4(%rsi,%rcx,4), %xmm1
	comiss	%xmm0, %xmm1
	jb	.L14
	cvtsi2ssl	%eax, %xmm0
	ret
.L14:
	movq	%rcx, %rax
	jmp	.L8
.L16:
	decl	%edx
	cvtsi2ssl	%edx, %xmm0
	ret
	.cfi_endproc
.LFE4001:
	.size	_Z13sampleFromCdfPjPfi, .-_Z13sampleFromCdfPjPfi
	.type	_Z13sampleFromPdfPjPfi.part.0, @function
_Z13sampleFromPdfPjPfi.part.0:
.LFB4656:
	.cfi_startproc
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movss	(%rsi), %xmm0
	movl	$1, %ecx
	xorps	%xmm1, %xmm1
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movss	%xmm0, 8(%rsp)
	movq	%rax, %r8
.L18:
	cmpl	%ecx, %edx
	jle	.L25
	movss	4(%rax), %xmm2
	movss	(%rax), %xmm0
	comiss	%xmm1, %xmm2
	jb	.L19
	addss	(%rsi,%rcx,4), %xmm0
.L19:
	movss	%xmm0, 4(%rax)
	incq	%rcx
	addq	$4, %rax
	jmp	.L18
.L25:
	movq	%r8, %rsi
	call	_Z13sampleFromCdfPjPfi
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L22
	call	__stack_chk_fail@PLT
.L22:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4656:
	.size	_Z13sampleFromPdfPjPfi.part.0, .-_Z13sampleFromPdfPjPfi.part.0
	.globl	_Z13sampleFromPdfPjPfi
	.type	_Z13sampleFromPdfPjPfi, @function
_Z13sampleFromPdfPjPfi:
.LFB4002:
	.cfi_startproc
	endbr64
	testl	%edx, %edx
	je	.L27
	jmp	_Z13sampleFromPdfPjPfi.part.0
.L27:
	xorps	%xmm0, %xmm0
	ret
	.cfi_endproc
.LFE4002:
	.size	_Z13sampleFromPdfPjPfi, .-_Z13sampleFromPdfPjPfi
	.globl	_Z3sumPfi
	.type	_Z3sumPfi, @function
_Z3sumPfi:
.LFB4003:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L29:
	cmpl	%eax, %esi
	jle	.L31
	addss	(%rdi,%rax,4), %xmm0
	incq	%rax
	jmp	.L29
.L31:
	ret
	.cfi_endproc
.LFE4003:
	.size	_Z3sumPfi, .-_Z3sumPfi
	.globl	_Z12swapZoneTypeP9zone_typeS0_
	.type	_Z12swapZoneTypeP9zone_typeS0_, @function
_Z12swapZoneTypeP9zone_typeS0_:
.LFB4004:
	.cfi_startproc
	endbr64
	movl	(%rsi), %edx
	movq	(%rdi), %rax
	movl	%edx, (%rdi)
	movl	4(%rsi), %edx
	movl	%edx, 4(%rdi)
	movq	%rax, (%rsi)
	ret
	.cfi_endproc
.LFE4004:
	.size	_Z12swapZoneTypeP9zone_typeS0_, .-_Z12swapZoneTypeP9zone_typeS0_
	.globl	_Z4dot2PfS_
	.type	_Z4dot2PfS_, @function
_Z4dot2PfS_:
.LFB4005:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L34:
	movss	(%rdi,%rax), %xmm1
	mulss	(%rsi,%rax), %xmm1
	addq	$4, %rax
	addss	%xmm1, %xmm0
	cmpq	$36, %rax
	jne	.L34
	ret
	.cfi_endproc
.LFE4005:
	.size	_Z4dot2PfS_, .-_Z4dot2PfS_
	.globl	_Z5noisefff
	.type	_Z5noisefff, @function
_Z5noisefff:
.LFB4006:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movaps	%xmm0, %xmm3
	movss	.LC5(%rip), %xmm0
	movss	%xmm2, 12(%rsp)
	movss	.LC4(%rip), %xmm2
	subss	%xmm3, %xmm0
	subss	%xmm1, %xmm2
	movaps	%xmm2, %xmm1
	call	fmaxf@PLT
	movss	.LC6(%rip), %xmm2
	subss	12(%rsp), %xmm2
	movaps	%xmm0, %xmm1
	movaps	%xmm2, %xmm0
	call	fmaxf@PLT
	xorps	%xmm1, %xmm1
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	jmp	fmaxf@PLT
	.cfi_endproc
.LFE4006:
	.size	_Z5noisefff, .-_Z5noisefff
	.globl	_Z9pollutionf
	.type	_Z9pollutionf, @function
_Z9pollutionf:
.LFB4007:
	.cfi_startproc
	endbr64
	movaps	%xmm0, %xmm1
	movss	.LC5(%rip), %xmm0
	subss	%xmm1, %xmm0
	xorps	%xmm1, %xmm1
	jmp	fmaxf@PLT
	.cfi_endproc
.LFE4007:
	.size	_Z9pollutionf, .-_Z9pollutionf
	.section	.rodata
	.align 32
.LC0:
	.long	1032358472
	.long	1032358472
	.long	1032358472
	.long	1045891645
	.long	1035489772
	.long	1035489772
	.long	1035489772
	.long	1039516303
	.long	1036831949
	.long	1036831949
	.align 32
.LC1:
	.long	1045220557
	.long	1052938076
	.long	1045220557
	.long	1031127695
	.long	1028443341
	.long	1022739087
	.long	1017370378
	.long	1008981770
	.long	1008981770
	.long	1017370378
	.long	0
	.long	0
	.long	1008981770
	.long	0
	.long	0
	.long	1008981770
	.long	0
	.long	0
	.text
	.globl	_Z8MCMCstepijP9zone_planS0_S0_
	.type	_Z8MCMCstepijP9zone_planS0_S0_, @function
_Z8MCMCstepijP9zone_planS0_S0_:
.LFB4008:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
	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
	movq	%rcx, %rbp
	xorl	%ecx, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdx, %rbx
	xorl	%edx, %edx
	subq	$648, %rsp
	.cfi_def_cfa_offset 704
	movl	%edi, 44(%rsp)
	xorl	%edi, %edi
	movl	%esi, 60(%rsp)
	movl	$520932557, %esi
	movq	%r8, 16(%rsp)
	salq	$33, %rsi
	xorl	%r8d, %r8d
	movq	%fs:40, %rax
	movq	%rax, 632(%rsp)
	xorl	%eax, %eax
	movq	%rdx, 272(%rsp)
	movl	$1045220557, %edx
	movabsq	$4474776592515176858, %rax
	movq	%rax, 280(%rsp)
	salq	$32, %rdx
	movabsq	$4453159313439771853, %rax
	movq	%rax, 296(%rsp)
	movl	$1057803469, %eax
	salq	$32, %rax
	movq	%rsi, 312(%rsp)
	movq	%rax, 320(%rsp)
	movabsq	$4453159313451515904, %rax
	movq	%rax, 368(%rsp)
	movabsq	$4483423504160102810, %rax
	movq	%rdx, 328(%rsp)
	movq	%rcx, 344(%rsp)
	movabsq	$4489188110458735821, %rcx
	movq	%rdi, 360(%rsp)
	movabsq	$4453159313448160461, %rdi
	movq	$1050253722, 288(%rsp)
	movq	$1045220557, 304(%rsp)
	movq	$1036831949, 336(%rsp)
	movq	$1028443341, 352(%rsp)
	movq	%rax, 376(%rsp)
	movabsq	$4489188110456051466, %rax
	movq	%rax, 400(%rsp)
	movabsq	$4453159313453193626, %rax
	movq	%rax, 424(%rsp)
	movl	$1028443341, %eax
	salq	$32, %rax
	movq	%rdi, 464(%rsp)
	movabsq	$4510805389534140826, %rdi
	movq	%rax, 448(%rsp)
	movl	$1036831949, %eax
	salq	$32, %rax
	movq	%rdi, 480(%rsp)
	movabsq	$4453159313444805018, %rdi
	movq	%rax, 456(%rsp)
	movq	%rax, 472(%rsp)
	movq	%rax, 504(%rsp)
	movq	%rax, 544(%rsp)
	movabsq	$4417130516425841050, %rax
	movq	%rdi, 488(%rsp)
	movabsq	$4489188110467124429, %rdi
	movq	%rax, 568(%rsp)
	movabsq	$4503599628419072000, %rax
	movq	%rdi, 520(%rsp)
	movl	$125, %edi
	movq	%rax, 584(%rsp)
	movl	$1053609165, %eax
	salq	$55, %rdi
	movq	%rdx, 600(%rsp)
	salq	$32, %rax
	movq	%rdx, 616(%rsp)
	leaq	128(%rsp), %rdx
	movq	%rcx, 440(%rsp)
	movq	%rsi, 496(%rsp)
	movq	%rcx, 512(%rsp)
	movq	%rdi, 528(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 536(%rsp)
	leaq	.LC0(%rip), %rsi
	movq	%rcx, 552(%rsp)
	movl	$10, %ecx
	movq	%rax, 592(%rsp)
	movabsq	$4656722015766446080, %rax
	movq	$1043207291, 384(%rsp)
	movq	$1043207291, 392(%rsp)
	movq	$1045220557, 408(%rsp)
	movq	$1050253722, 416(%rsp)
	movq	%r8, 432(%rsp)
	movq	$1050253722, 560(%rsp)
	movq	%r9, 576(%rsp)
	movq	%r10, 608(%rsp)
	movq	$1045220557, 624(%rsp)
	rep movsl
	leaq	128(%rsp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$18, %ecx
	rep movsl
	movl	$18, %esi
	movq	%rdx, %rdi
	movq	%rax, 76(%rsp)
	movl	$0x41200000, 84(%rsp)
	call	_Z3sumPfi
	leaq	200(%rsp), %rsi
	xorl	%eax, %eax
	movq	%rsi, 24(%rsp)
.L40:
	movss	(%rdx,%rax), %xmm1
	movq	24(%rsp), %rsi
	addq	$4, %rax
	divss	%xmm0, %xmm1
	mulss	.LC8(%rip), %xmm1
	mulss	.LC8(%rip), %xmm1
	movss	%xmm1, -4(%rax,%rsi)
	cmpq	$72, %rax
	jne	.L40
	movq	%rbx, %r13
	leaq	800(%rbx), %r15
.L41:
	movq	%r13, %r12
	movl	$10, %r14d
.L42:
	leaq	60(%rsp), %rax
	movq	24(%rsp), %rsi
	movl	$18, %edx
	addq	$8, %r12
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	_Z13sampleFromPdfPjPfi.part.0
	movl	$3, %esi
	cvttss2sil	%xmm0, %ecx
	movl	%ecx, %eax
	movslq	%ecx, %rcx
	movss	200(%rsp,%rcx,4), %xmm0
	cltd
	subss	.LC9(%rip), %xmm0
	idivl	%esi
	movss	%xmm0, 200(%rsp,%rcx,4)
	incl	%edx
	movl	%eax, -8(%r12)
	movl	%edx, -4(%r12)
	decl	%r14d
	jne	.L42
	addq	$80, %r13
	cmpq	%r15, %r13
	jne	.L41
	movq	16(%rsp), %rax
	leaq	.L52(%rip), %r15
	movl	$0x00000000, 800(%rbx)
	movl	$0x00000000, 800(%rax)
	xorl	%eax, %eax
	movl	%eax, 12(%rsp)
.L44:
	movl	44(%rsp), %ecx
	cmpl	%ecx, 12(%rsp)
	jge	.L84
	movl	$201, %ecx
	movq	%rbp, %rdi
	xorps	%xmm1, %xmm1
	movq	%rbx, %rsi
	rep movsl
.L46:
	movq	(%rsp), %rdi
	call	_Z5randfPj
	mulss	.LC8(%rip), %xmm0
	movq	(%rsp), %rdi
	addss	%xmm1, %xmm0
	cvttss2sil	%xmm0, %ecx
	call	_Z5randfPj
	mulss	.LC8(%rip), %xmm0
	movq	(%rsp), %rdi
	movslq	%ecx, %rcx
	addss	%xmm1, %xmm0
	cvttss2sil	%xmm0, %edx
	call	_Z5randfPj
	mulss	.LC8(%rip), %xmm0
	movq	(%rsp), %rdi
	movslq	%edx, %rdx
	imulq	$10, %rdx, %r8
	addss	%xmm1, %xmm0
	addq	%rcx, %r8
	cvttss2sil	%xmm0, %esi
	call	_Z5randfPj
	mulss	.LC8(%rip), %xmm0
	movslq	%esi, %rsi
	addss	%xmm1, %xmm0
	cvttss2sil	%xmm0, %eax
	cltq
	imulq	$10, %rax, %rdi
	addq	%rsi, %rdi
	movl	0(%rbp,%rdi,8), %r9d
	cmpl	%r9d, 0(%rbp,%r8,8)
	jne	.L45
	movl	4(%rbp,%rdi,8), %edi
	cmpl	%edi, 4(%rbp,%r8,8)
	je	.L46
.L45:
	imulq	$10, %rax, %rax
	xorl	%r13d, %r13d
	imulq	$10, %rdx, %rdx
	addq	%rsi, %rax
	addq	%rcx, %rdx
	leaq	0(%rbp,%rax,8), %rsi
	leaq	0(%rbp,%rdx,8), %rdi
	call	_Z12swapZoneTypeP9zone_typeS0_
	movl	$0x00000000, 32(%rsp)
	movl	$0x00000000, 800(%rbp)
.L47:
	imulq	$80, %r13, %r12
	xorl	%r14d, %r14d
.L60:
	movl	(%rbx,%r12), %ecx
	testl	%ecx, %ecx
	jne	.L48
	movss	.LC7(%rip), %xmm3
	movq	%rbp, %r8
	movaps	%xmm3, %xmm1
	movaps	%xmm3, %xmm4
	movss	%xmm3, 8(%rsp)
	movaps	%xmm3, %xmm2
.L49:
	movl	%r13d, %edx
	xorl	%esi, %esi
	subl	%ecx, %edx
	movl	%edx, %edi
	negl	%edi
	cmovs	%edx, %edi
.L57:
	movl	(%r8,%rsi,8), %edx
	testl	%edx, %edx
	je	.L50
	movl	%r14d, %r11d
	subl	%esi, %r11d
	movl	%r11d, %r10d
	negl	%r10d
	cmovs	%r11d, %r10d
	decl	%edx
	addl	%edi, %r10d
	imull	$400, %r10d, %r10d
	cvtsi2ssl	%r10d, %xmm0
	cmpl	$4, %edx
	ja	.L50
	movslq	(%r15,%rdx,4), %rdx
	addq	%r15, %rdx
	notrack jmp	*%rdx
	.section	.rodata
	.align 4
	.align 4
.L52:
	.long	.L56-.L52
	.long	.L55-.L52
	.long	.L54-.L52
	.long	.L53-.L52
	.long	.L51-.L52
	.text
.L56:
	minss	%xmm2, %xmm0
	movaps	%xmm0, %xmm2
	jmp	.L50
.L55:
	minss	8(%rsp), %xmm0
	movss	%xmm0, 8(%rsp)
	jmp	.L50
.L54:
	minss	%xmm4, %xmm0
	movaps	%xmm0, %xmm4
	jmp	.L50
.L53:
	minss	%xmm1, %xmm0
	movaps	%xmm0, %xmm1
	jmp	.L50
.L51:
	minss	%xmm3, %xmm0
	movaps	%xmm0, %xmm3
.L50:
	incq	%rsi
	cmpq	$10, %rsi
	jne	.L57
	incl	%ecx
	addq	$80, %r8
	cmpl	$10, %ecx
	jne	.L49
	movss	.LC10(%rip), %xmm5
	movss	%xmm2, 48(%rsp)
	movss	%xmm1, 56(%rsp)
	mulss	%xmm2, %xmm5
	movss	%xmm4, 52(%rsp)
	movss	%xmm3, 40(%rsp)
	movaps	%xmm5, %xmm0
	movss	%xmm5, 36(%rsp)
	call	expf@PLT
	movss	40(%rsp), %xmm3
	mulss	.LC10(%rip), %xmm3
	movss	%xmm0, 200(%rsp)
	movaps	%xmm3, %xmm0
	movss	%xmm3, 40(%rsp)
	call	expf@PLT
	movss	%xmm0, 204(%rsp)
	movss	36(%rsp), %xmm0
	call	expf@PLT
	movss	52(%rsp), %xmm4
	movss	%xmm0, 208(%rsp)
	movss	.LC10(%rip), %xmm0
	mulss	%xmm4, %xmm0
	call	expf@PLT
	movss	56(%rsp), %xmm1
	movss	%xmm0, 212(%rsp)
	movss	.LC10(%rip), %xmm0
	movss	%xmm1, 36(%rsp)
	mulss	%xmm1, %xmm0
	call	expf@PLT
	movss	%xmm0, 216(%rsp)
	movss	40(%rsp), %xmm0
	call	expf@PLT
	movss	48(%rsp), %xmm2
	movss	36(%rsp), %xmm1
	movss	%xmm0, 220(%rsp)
	movss	8(%rsp), %xmm0
	call	_Z5noisefff
	mulss	.LC10(%rip), %xmm0
	call	expf@PLT
	movss	%xmm0, 224(%rsp)
	movss	8(%rsp), %xmm0
	call	_Z9pollutionf
	mulss	.LC10(%rip), %xmm0
	call	expf@PLT
	movl	4(%rbp,%r12), %eax
	movl	$0x00000000, 232(%rsp)
	leaq	272(%rsp), %rcx
	movss	%xmm0, 228(%rsp)
	leal	-1(%rax), %edx
	movslq	%edx, %rdx
	movss	76(%rsp,%rdx,4), %xmm2
	xorl	%edx, %edx
.L59:
	imulq	$36, %rdx, %rdi
	movq	24(%rsp), %rsi
	addq	%rcx, %rdi
	call	_Z4dot2PfS_
	movss	88(%rsp,%rdx,4), %xmm1
	incq	%rdx
	mulss	%xmm1, %xmm0
	mulss	%xmm2, %xmm1
	addss	32(%rsp), %xmm1
	mulss	%xmm2, %xmm0
	movss	%xmm1, 32(%rsp)
	addss	800(%rbp), %xmm0
	movss	%xmm0, 800(%rbp)
	cmpq	$10, %rdx
	jne	.L59
.L48:
	incl	%r14d
	addq	$8, %r12
	cmpl	$10, %r14d
	jne	.L60
	incq	%r13
	cmpq	$10, %r13
	jne	.L47
	movq	16(%rsp), %rax
	movss	800(%rbp), %xmm0
	divss	32(%rsp), %xmm0
	movss	%xmm0, 800(%rbp)
	comiss	800(%rax), %xmm0
	jbe	.L62
	movl	$201, %ecx
	movq	%rax, %rdi
	movq	%rbp, %rsi
	rep movsl
.L62:
	movss	800(%rbp), %xmm0
	comiss	800(%rbx), %xmm0
	ja	.L64
	movl	12(%rsp), %eax
	movl	$10, %ecx
	cltd
	idivl	%ecx
	testl	%edx, %edx
	jne	.L65
.L64:
	movl	$201, %ecx
	movq	%rbx, %rdi
	movq	%rbp, %rsi
	rep movsl
.L65:
	incl	12(%rsp)
	jmp	.L44
.L84:
	movq	632(%rsp), %rax
	subq	%fs:40, %rax
	je	.L67
	call	__stack_chk_fail@PLT
.L67:
	addq	$648, %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
.LFE4008:
	.size	_Z8MCMCstepijP9zone_planS0_S0_, .-_Z8MCMCstepijP9zone_planS0_S0_
	.globl	_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_
	.type	_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_, @function
_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_:
.LFB4034:
	.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
	movq	%rcx, (%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)
	movq	%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	.L85
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_(%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
.L85:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L87
	call	__stack_chk_fail@PLT
.L87:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4034:
	.size	_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_, .-_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_
	.globl	_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_
	.type	_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_, @function
_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_:
.LFB4035:
	.cfi_startproc
	endbr64
	jmp	_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_
	.cfi_endproc
.LFE4035:
	.size	_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_, .-_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC13:
	.string	"cuda memory allocation error!\n"
.LC14:
	.string	"cuda memory copy error!\n"
.LC15:
	.string	"start GPU kernel.\n"
.LC16:
	.string	"GPU kernel done.\n"
	.text
	.globl	_Z19zonePlanMCMCGPUfuncPP9zone_plani
	.type	_Z19zonePlanMCMCGPUfuncPP9zone_plani, @function
_Z19zonePlanMCMCGPUfuncPP9zone_plani:
.LFB4009:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	movl	$804, %edi
	subq	$80, %rsp
	.cfi_def_cfa_offset 96
	movl	%esi, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	leaq	16(%rsp), %rdi
	movl	$4, %esi
	movq	%rax, (%rbx)
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L104
	leaq	24(%rsp), %rdi
	movl	$804, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L104
	leaq	32(%rsp), %rdi
	movl	$804, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L94
	movq	24(%rsp), %rdi
	jmp	.L105
.L94:
	leaq	40(%rsp), %rdi
	movl	$804, %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L95
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
.L105:
	call	cudaFree@PLT
.L104:
	leaq	.LC13(%rip), %rsi
	jmp	.L102
.L95:
	movq	16(%rsp), %rdi
	leaq	12(%rsp), %rsi
	movl	$1, %ecx
	movl	$4, %edx
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L96
	movq	24(%rsp), %rdi
	jmp	.L103
.L96:
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rdi, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L97
	movq	40(%rsp), %rcx
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	call	_Z58__device_stub__Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_PiP9zone_planS1_S1_
.L97:
	movq	(%rbx), %rdi
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movl	$804, %edx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	testl	%eax, %eax
	je	.L98
.L103:
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	leaq	.LC14(%rip), %rsi
	jmp	.L102
.L98:
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	leaq	.LC16(%rip), %rsi
.L102:
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L100
	call	__stack_chk_fail@PLT
.L100:
	addq	$80, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4009:
	.size	_Z19zonePlanMCMCGPUfuncPP9zone_plani, .-_Z19zonePlanMCMCGPUfuncPP9zone_plani
	.section	.rodata.str1.1
.LC17:
	.string	"_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4037:
	.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	.LC17(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z21zonePlanMCMCGPUKernelPiP9zone_planS1_S1_(%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
.LFE4037:
	.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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC2:
	.long	805306368
	.align 4
.LC4:
	.long	1137180672
	.align 4
.LC5:
	.long	1145569280
	.align 4
.LC6:
	.long	1128792064
	.align 4
.LC7:
	.long	1165623296
	.align 4
.LC8:
	.long	1092616192
	.align 4
.LC9:
	.long	1065353216
	.align 4
.LC10:
	.long	-1165815185
	.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:

#ifndef __ZONE_PLAN_MCMC_CUH__

#include <stdio.h>

#define ZONE_PLAN_MCMC_GRID_SIZE	1 // 8
#define ZONE_PLAN_MCMC_BLOCK_SIZE	1 // 128

#define ZONE_GRID_SIZE 10 // 200
#define ZONE_CELL_LEN 400 // 20

struct zone_type {
	int type;
	int level;
};

struct zone_plan {
	zone_type zones[ZONE_GRID_SIZE][ZONE_GRID_SIZE];
	float score;
};

void MCMCstep(int numIterations, unsigned int randx, zone_plan* plan, zone_plan* proposal, zone_plan* bestPlan);

void zonePlanMCMCGPUfunc(zone_plan** bestPlan, int numIterations);

#endif

	.text
	.file	"ZonePlanMCMC.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
