/******************************************************************************
LICENSE

Copyright (c) 2015 Codeplay Software Ltd.
Copyright (c) 2006-2008 Kevin Beason (kevin.beason@gmail.com)

Permission is hereby granted, free of charge, to any person obtaining
a copy of this software and associated documentation files (the
"Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:

The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 *
 * Modified version for SYCL of Kevin Beason smallpt
 * http://www.kevinbeason.com/smallpt/
 *
 *
 *
 * Modified version using CUDA
 *   Add the original License in the source file
 *   The number of samples is set at 40
 *   spheres_glob is not declared in the global space
 *
 * Zheming Jin
 *
 ******************************************************************************/

#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <cuda.h>


class RNG {
	public:
		unsigned int x;
		const uint32_t fmask = (1 << 23) - 1;

	        __device__
		RNG(const unsigned int seed) { x = seed; }

	        __device__
		uint32_t next() {
			x ^= x >> 6;
			x ^= x << 17;
			x ^= x >> 9;
			return uint32_t(x);
		}
	        __device__
		float operator()(void) {
			union {
				float f;
				uint32_t i;
			} u;
			u.i = (next() & fmask) | 0x3f800000;
			return u.f - 1.f;
		}
};

struct Vec {     // Usage: time ./smallpt 5000 && xv image.ppm
	float x, y, z; // position, also color (r,g,b)
	__host__ __device__
	Vec(float x_ = 0, float y_ = 0, float z_ = 0) : x(x_), y(y_), z(z_) {}
	__device__
	Vec operator+(const Vec &b) const { return Vec(x + b.x, y + b.y, z + b.z); }
	__device__
	Vec operator-(const Vec &b) const { return Vec(x - b.x, y - b.y, z - b.z); }
	__host__ __device__
	Vec operator*(float b) const { return Vec(x * b, y * b, z * b); }
	__device__
	Vec mult(const Vec &b) const { return Vec(x * b.x, y * b.y, z * b.z); }
	__device__
	Vec &norm() {
		return *this = *this * (1 / sqrt(x * x + y * y + z * z));
	}
	__device__
	float dot(const Vec &b) const {
		return x * b.x + y * b.y + z * b.z;
	} // cross:
	__device__
	Vec operator%(Vec &b) {
		return Vec(y * b.z - z * b.y, z * b.x - x * b.z, x * b.y - y * b.x);
	}
};

struct Ray {
	Vec o, d;
        __device__
	Ray(Vec o_, Vec d_) : o(o_), d(d_) {}
};

enum Refl_t { DIFF, SPEC, REFR }; // material types, used in radiance()

struct Sphere {
	float rad;   // radius
	Vec p, e, c; // position, emission, color
	Refl_t refl; // reflection type (DIFFuse, SPECular, REFRactive)
	__host__ __device__
	Sphere(float rad_, Vec p_, Vec e_, Vec c_, Refl_t refl_)
		: rad(rad_), p(p_), e(e_), c(c_), refl(refl_) {}
	__device__
	inline float intersect(const Ray &r) const { // returns distance, 0 if nohit
		Vec op = p - r.o; // Solve t^2*d.d + 2*t*(o-p).d + (o-p).(o-p)-R^2 = 0
		float t, eps = 1.5e-2f, b = op.dot(r.d),
		      det = b * b - op.dot(op) + rad * rad;
		if (det < 0)
			return 0;
		else
			det = sqrt(det);
		return (t = b - det) > eps ? t : ((t = b + det) > eps ? t : 0);
	}
};

__host__ __device__
inline float clamp(float x) { return x < 0 ? 0 : x > 1 ? 1 : x; }

inline int toInt(float x) {
	return int(pow(clamp(x), 1 / 2.2f) * 255 + .5f);
}

template<typename T>
__device__
inline bool intersect(const Ray &r, float &t, int &id,
		T spheres) {
	float d, inf = t = 1e20f;
	for (int i = 9; i--;)
		if ((d = spheres[i].intersect(r)) && d < t) {
			t = d;
			id = i;
		}
	return t < inf;
}

template<typename T>
__device__
Vec radiance(const Ray &r_, int depth_, T spheres, RNG &rng) {
	float t;
	int id = 0;
	Ray r = r_;
	int depth = depth_;
	Vec cl(0, 0, 0); // accumulated color
	Vec cf(1, 1, 1); // accumulated reflectance
	while (1) {
		if (!intersect(r, t, id, spheres))
			return cl;                     // if miss, return black
		const Sphere &obj = spheres[id]; // the hit object
		Vec x = r.o + r.d * t, n = (x - obj.p).norm(),
		    nl = n.dot(r.d) < 0 ? n : n * -1, f = obj.c;
		float p = f.x > f.y && f.x > f.z ? f.x : f.y > f.z ? f.y : f.z; // max refl
		cl = cl + cf.mult(obj.e);
		if (++depth > 5) {
			if (rng() < p) {
				f = f * (1 / p);
			} else {
				return cl;
			}
		} // R.R.
		cf = cf.mult(f);
		if (obj.refl == DIFF) { // Ideal DIFFUSE reflection
			float r1 = 2 * M_PI * rng(), r2 = rng(), r2s = sqrt(r2);
			Vec w = nl,
			    u = ((fabs(w.x) > .1 ? Vec(0, 1) : Vec(1)) % w).norm(),
			    v = w % u;
			Vec d = (u * cos(r1) * r2s + v * sin(r1) * r2s +
					w * sqrt(1 - r2)).norm();
			r = Ray(x, d);
			continue;
		} else if (obj.refl == SPEC) { // Ideal SPECULAR reflection
			r = Ray(x, r.d - n * 2 * n.dot(r.d));
			continue;
		}
		Ray reflRay(x, r.d - n * 2 * n.dot(r.d)); // Ideal dielectric REFRACTION
		bool into = n.dot(nl) > 0;                // Ray from outside going in?
		float nc = 1, nt = 1.5, nnt = into ? nc / nt : nt / nc, ddn = r.d.dot(nl),
		      cos2t;
		if ((cos2t = 1 - nnt * nnt * (1 - ddn * ddn)) <
				0) { // Total internal reflection
			r = reflRay;
			continue;
		}
		Vec tdir =
			(r.d * nnt -
			 n * ((into ? 1 : -1) * (ddn * nnt + sqrt(cos2t)))).norm();
		float a = nt - nc, b = nt + nc, R0 = a * a / (b * b),
		      c = 1 - (into ? -ddn : tdir.dot(n));
		float Re = R0 + (1 - R0) * c * c * c * c * c, Tr = 1 - Re,
		      P = .25 + .5 * Re, RP = Re / P, TP = Tr / (1 - P);
		if (rng() < P) {
			cf = cf * RP;
			r = reflRay;
		} else {
			cf = cf * TP;
			r = Ray(x, tdir);
		}
		continue;
	}
}

__global__ void raytrace(Vec* c, const Sphere* spheres_, const int w, const int h, const int samps) {
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;
	Vec r;
	const Sphere *spheres = &spheres_[0];
	Ray cam(Vec(50, 52, 295.6), Vec(0, -0.042612, -1).norm()); // cam pos, dir
	Vec cx = Vec(w * .5135 / h), cy = (cx % cam.d).norm() * .5135;
	RNG rng(1 + (y * w) + x); // initialise our own rng with rand() seed
	for (int sy = 0, i = (h - y - 1) * w + x; sy < 2; sy++) // 2x2 subpixel rows
		for (int sx = 0; sx < 2; sx++, r = Vec()) {           // 2x2 subpixel cols
			for (int s = 0; s < samps; s++) {
				float r1 = 2 * rng(), dx = r1 < 1 ? sqrt(r1) - 1
					: 1 - sqrt(2 - r1);
				float r2 = 2 * rng(), dy = r2 < 1 ? sqrt(r2) - 1
					: 1 - sqrt(2 - r2);
				Vec d = cx * (((sx + .5 + dx) / 2 + x) / w - .5) +
					cy * (((sy + .5 + dy) / 2 + y) / h - .5) + cam.d;
				r = r + radiance(Ray(cam.o + d * 140, d.norm()), 0, spheres, rng) *
					(1. / samps);
			} // Camera rays are pushed ^^^^^ forward to start in interior
			c[i] = c[i] + Vec(clamp(r.x), clamp(r.y), clamp(r.z)) * .25;
		}
};

int main(int argc, char *argv[]) {

	const Sphere spheres_glob[] = {
		// Scene: radius, position, emission, color, material
		Sphere(1e4, Vec(1e4 + 1, 40.8, 81.6), Vec(), Vec(.75, .25, .25),
				DIFF), // Left
		Sphere(1e4, Vec(-1e4 + 99, 40.8, 81.6), Vec(), Vec(.25, .25, .75),
				DIFF),                                                     // Rght
		Sphere(1e4, Vec(50, 40.8, 1e4), Vec(), Vec(.75, .75, .75), DIFF), // Back
		Sphere(1e4, Vec(50, 40.8, -1e4 + 170), Vec(), Vec(), DIFF),       // Frnt
		Sphere(1e4, Vec(50, 1e4, 81.6), Vec(), Vec(.75, .75, .75), DIFF), // Botm
		Sphere(1e4, Vec(50, -1e4 + 81.6, 81.6), Vec(), Vec(.75, .75, .75),
				DIFF),                                                      // Top
		Sphere(16.5, Vec(27, 16.5, 47), Vec(), Vec(1, 1, 1) * .999, SPEC), // Mirr
		Sphere(16.5, Vec(73, 16.5, 78), Vec(), Vec(1, 1, 1) * .999, REFR), // Glas
		Sphere(600, Vec(50, 681.6 - .27, 81.6), Vec(12, 12, 12), Vec(),
				DIFF) // Lite
	};

	int w = 1024, h = 768; 
	int samps = 40; // # samples
	Vec *c = (Vec*) malloc(sizeof(Vec) * w * h);

	Vec *color_buffer;
	Sphere *spheres_buffer;

	cudaMalloc((void**)&color_buffer, sizeof(Vec) * w * h);
	cudaMalloc((void**)&spheres_buffer, sizeof(Sphere) * 9);

	cudaMemcpy(spheres_buffer, &spheres_glob[0], sizeof(Sphere) * 9, cudaMemcpyHostToDevice);

	raytrace<<< dim3(w/8, h/8), dim3(8, 8) >>>(color_buffer, spheres_buffer, w, h, samps);

	cudaMemcpy(c, color_buffer, sizeof(Vec) * w * h, cudaMemcpyDeviceToHost);

	FILE *f = fopen("image.ppm", "w"); // Write image to PPM file.
	fprintf(f, "P3\n%d %d\n%d\n", w, h, 255);
	for (int i = 0; i < w * h; i++)
		fprintf(f, "%d %d %d ", toInt(c[i].x), toInt(c[i].y), toInt(c[i].z));

	free(c);
	cudaFree(color_buffer);
	cudaFree(spheres_buffer);
	return 0;
}

	.file	"tmpxft_00329a5c_00000000-6_main.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2056:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2056:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text._Z5toIntf,"axG",@progbits,_Z5toIntf,comdat
	.weak	_Z5toIntf
	.type	_Z5toIntf, @function
_Z5toIntf:
.LFB2050:
	.cfi_startproc
	endbr64
	xorps	%xmm1, %xmm1
	pushq	%rcx
	.cfi_def_cfa_offset 16
	comiss	%xmm0, %xmm1
	ja	.L5
	movss	.LC1(%rip), %xmm1
	minss	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	jmp	.L3
.L5:
	xorps	%xmm0, %xmm0
.L3:
	movss	.LC2(%rip), %xmm1
	call	powf@PLT
	mulss	.LC3(%rip), %xmm0
	addss	.LC4(%rip), %xmm0
	popq	%rdx
	.cfi_def_cfa_offset 8
	cvttss2sil	%xmm0, %eax
	ret
	.cfi_endproc
.LFE2050:
	.size	_Z5toIntf, .-_Z5toIntf
	.text
	.globl	_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii
	.type	_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii, @function
_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii:
.LFB2078:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%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, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z8raytraceP3VecPK6Sphereiii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L9:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2078:
	.size	_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii, .-_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii
	.globl	_Z8raytraceP3VecPK6Sphereiii
	.type	_Z8raytraceP3VecPK6Sphereiii, @function
_Z8raytraceP3VecPK6Sphereiii:
.LFB2079:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii
	.cfi_endproc
.LFE2079:
	.size	_Z8raytraceP3VecPK6Sphereiii, .-_Z8raytraceP3VecPK6Sphereiii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC5:
	.string	"w"
.LC6:
	.string	"image.ppm"
.LC7:
	.string	"P3\n%d %d\n%d\n"
.LC8:
	.string	"%d %d %d "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2053:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	xorl	%esi, %esi
	movl	$253, %ecx
	movl	$71793, %edx
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	salq	$54, %rcx
	salq	$46, %rdx
	xorl	%edi, %edi
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	xorl	%r11d, %r11d
	xorl	%ebp, %ebp
	xorl	%r12d, %r12d
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movabsq	$4765709128148104192, %rbx
	subq	$464, %rsp
	.cfi_def_cfa_offset 512
	movq	%fs:40, %rax
	movq	%rax, 456(%rsp)
	movabsq	$5051987649969799168, %rax
	movq	%rbx, 108(%rsp)
	movabsq	$4765709125936480256, %rbx
	movq	%rsi, 124(%rsp)
	movabsq	$4776067406002667520, %rsi
	movq	%rbx, 196(%rsp)
	movl	$3323566080, %ebx
	movq	%rsi, 148(%rsp)
	movq	%rcx, 84(%rsp)
	movq	%rdx, 100(%rsp)
	movq	%rax, 60(%rsp)
	movabsq	$4801737922953032499, %rax
	movq	%rax, 68(%rsp)
	xorl	%eax, %eax
	movq	%rax, 76(%rsp)
	movabsq	$4503599628419072000, %rax
	movq	%rax, 92(%rsp)
	movq	%rax, 132(%rsp)
	movabsq	$5051983251856634675, %rax
	movq	%rax, 156(%rsp)
	movabsq	$4557642823960100864, %rax
	movq	%rdi, 164(%rsp)
	movl	$9437184, %edi
	movq	%rcx, 172(%rsp)
	movq	%rax, 180(%rsp)
	movq	%rdx, 188(%rsp)
	movq	$1117991731, 116(%rsp)
	movq	$1061158912, 140(%rsp)
	movq	%rbx, 204(%rsp)
	movabsq	$4801737923019685888, %rbx
	movq	%rax, 268(%rsp)
	movq	%rax, 308(%rsp)
	movabsq	$4744542208533987328, %rax
	movq	%rax, 324(%rsp)
	movabsq	$4772689706205052928, %rax
	movq	%rax, 332(%rsp)
	movl	$1065336439, %eax
	salq	$32, %rax
	movq	%rbx, 244(%rsp)
	movabsq	$4720898309391122433, %rbx
	movq	%rax, 348(%rsp)
	addq	$1065336439, %rax
	movq	%rax, 356(%rsp)
	movq	%rax, 396(%rsp)
	movabsq	$9655271031, %rax
	movq	%rax, 404(%rsp)
	movabsq	$4776067405968703488, %rax
	movq	%rax, 412(%rsp)
	movabsq	$4801737922987054367, %rax
	movq	%rcx, 260(%rsp)
	movabsq	$-4171747663873048576, %rcx
	movq	%rdx, 276(%rsp)
	movabsq	$4720898310507986944, %rdx
	movq	%rbx, 364(%rsp)
	xorl	%ebx, %ebx
	movq	%rax, 420(%rsp)
	movabsq	$4701758012069511168, %rax
	movq	%r8, 212(%rsp)
	movq	%r8, 220(%rsp)
	movq	%r8, 228(%rsp)
	movq	%r9, 252(%rsp)
	movq	%rcx, 284(%rsp)
	movq	%r10, 300(%rsp)
	movq	%r11, 340(%rsp)
	movq	%rdx, 372(%rsp)
	movq	%rsi, 236(%rsp)
	movq	%rbx, 388(%rsp)
	movq	$1117991731, 292(%rsp)
	movq	$1061158912, 316(%rsp)
	movq	$1117519872, 380(%rsp)
	movq	%rax, 428(%rsp)
	movq	$1094713344, 436(%rsp)
	movq	%rbp, 444(%rsp)
	movl	%r12d, 452(%rsp)
	call	malloc@PLT
	leaq	16(%rsp), %rdi
	movl	$9437184, %esi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movl	$396, %esi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdi
	leaq	60(%rsp), %rsi
	movl	$1, %ecx
	movl	$396, %edx
	call	cudaMemcpy@PLT
	movl	$3221225473, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$7, %rdi
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$34359738376, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$40, %r8d
	movl	$768, %ecx
	movl	$1024, %edx
	call	_Z42__device_stub__Z8raytraceP3VecPK6SphereiiiP3VecPK6Sphereiii
.L15:
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movq	%rbx, %rbp
	movl	$9437184, %edx
	leaq	9437184(%rbx), %r14
	leaq	.LC8(%rip), %r13
	call	cudaMemcpy@PLT
	leaq	.LC5(%rip), %rsi
	leaq	.LC6(%rip), %rdi
	call	fopen@PLT
	movl	$255, %r9d
	movl	$768, %r8d
	movl	$1024, %ecx
	movq	%rax, %rdi
	movq	%rax, %r12
	movl	$2, %esi
	xorl	%eax, %eax
	leaq	.LC7(%rip), %rdx
	call	__fprintf_chk@PLT
.L16:
	movss	8(%rbp), %xmm0
	addq	$12, %rbp
	call	_Z5toIntf
	movss	-8(%rbp), %xmm0
	movl	%eax, 12(%rsp)
	call	_Z5toIntf
	movss	-12(%rbp), %xmm0
	movl	%eax, 8(%rsp)
	call	_Z5toIntf
	movl	12(%rsp), %r9d
	movq	%r13, %rdx
	movq	%r12, %rdi
	movl	8(%rsp), %r8d
	movl	%eax, %ecx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	cmpq	%r14, %rbp
	jne	.L16
	movq	%rbx, %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	456(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$464, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	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
.LFE2053:
	.size	main, .-main
	.section	.rodata.str1.1
.LC9:
	.string	"_Z8raytraceP3VecPK6Sphereiii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2081:
	.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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z8raytraceP3VecPK6Sphereiii(%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
.LFE2081:
	.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
.LC1:
	.long	1065353216
	.align 4
.LC2:
	.long	1055439406
	.align 4
.LC3:
	.long	1132396544
	.align 4
.LC4:
	.long	1056964608
	.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:

/******************************************************************************
LICENSE

Copyright (c) 2015 Codeplay Software Ltd.
Copyright (c) 2006-2008 Kevin Beason (kevin.beason@gmail.com)

Permission is hereby granted, free of charge, to any person obtaining
a copy of this software and associated documentation files (the
"Software"), to deal in the Software without restriction, including
without limitation the rights to use, copy, modify, merge, publish,
distribute, sublicense, and/or sell copies of the Software, and to
permit persons to whom the Software is furnished to do so, subject to
the following conditions:

The above copyright notice and this permission notice shall be included
in all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 *
 * Modified version for SYCL of Kevin Beason smallpt
 * http://www.kevinbeason.com/smallpt/
 *
 *
 *
 * Modified version using CUDA
 *   Add the original License in the source file
 *   The number of samples is set at 40
 *   spheres_glob is not declared in the global space
 *
 * Zheming Jin
 *
 ******************************************************************************/

#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <hip/hip_runtime.h>


class RNG {
	public:
		unsigned int x;
		const uint32_t fmask = (1 << 23) - 1;

	        __device__
		RNG(const unsigned int seed) { x = seed; }

	        __device__
		uint32_t next() {
			x ^= x >> 6;
			x ^= x << 17;
			x ^= x >> 9;
			return uint32_t(x);
		}
	        __device__
		float operator()(void) {
			union {
				float f;
				uint32_t i;
			} u;
			u.i = (next() & fmask) | 0x3f800000;
			return u.f - 1.f;
		}
};

struct Vec {     // Usage: time ./smallpt 5000 && xv image.ppm
	float x, y, z; // position, also color (r,g,b)
	__host__ __device__
	Vec(float x_ = 0, float y_ = 0, float z_ = 0) : x(x_), y(y_), z(z_) {}
	__device__
	Vec operator+(const Vec &b) const { return Vec(x + b.x, y + b.y, z + b.z); }
	__device__
	Vec operator-(const Vec &b) const { return Vec(x - b.x, y - b.y, z - b.z); }
	__host__ __device__
	Vec operator*(float b) const { return Vec(x * b, y * b, z * b); }
	__device__
	Vec mult(const Vec &b) const { return Vec(x * b.x, y * b.y, z * b.z); }
	__device__
	Vec &norm() {
		return *this = *this * (1 / sqrt(x * x + y * y + z * z));
	}
	__device__
	float dot(const Vec &b) const {
		return x * b.x + y * b.y + z * b.z;
	} // cross:
	__device__
	Vec operator%(Vec &b) {
		return Vec(y * b.z - z * b.y, z * b.x - x * b.z, x * b.y - y * b.x);
	}
};

struct Ray {
	Vec o, d;
        __device__
	Ray(Vec o_, Vec d_) : o(o_), d(d_) {}
};

enum Refl_t { DIFF, SPEC, REFR }; // material types, used in radiance()

struct Sphere {
	float rad;   // radius
	Vec p, e, c; // position, emission, color
	Refl_t refl; // reflection type (DIFFuse, SPECular, REFRactive)
	__host__ __device__
	Sphere(float rad_, Vec p_, Vec e_, Vec c_, Refl_t refl_)
		: rad(rad_), p(p_), e(e_), c(c_), refl(refl_) {}
	__device__
	inline float intersect(const Ray &r) const { // returns distance, 0 if nohit
		Vec op = p - r.o; // Solve t^2*d.d + 2*t*(o-p).d + (o-p).(o-p)-R^2 = 0
		float t, eps = 1.5e-2f, b = op.dot(r.d),
		      det = b * b - op.dot(op) + rad * rad;
		if (det < 0)
			return 0;
		else
			det = sqrt(det);
		return (t = b - det) > eps ? t : ((t = b + det) > eps ? t : 0);
	}
};

__host__ __device__
inline float clamp(float x) { return x < 0 ? 0 : x > 1 ? 1 : x; }

inline int toInt(float x) {
	return int(pow(clamp(x), 1 / 2.2f) * 255 + .5f);
}

template<typename T>
__device__
inline bool intersect(const Ray &r, float &t, int &id,
		T spheres) {
	float d, inf = t = 1e20f;
	for (int i = 9; i--;)
		if ((d = spheres[i].intersect(r)) && d < t) {
			t = d;
			id = i;
		}
	return t < inf;
}

template<typename T>
__device__
Vec radiance(const Ray &r_, int depth_, T spheres, RNG &rng) {
	float t;
	int id = 0;
	Ray r = r_;
	int depth = depth_;
	Vec cl(0, 0, 0); // accumulated color
	Vec cf(1, 1, 1); // accumulated reflectance
	while (1) {
		if (!intersect(r, t, id, spheres))
			return cl;                     // if miss, return black
		const Sphere &obj = spheres[id]; // the hit object
		Vec x = r.o + r.d * t, n = (x - obj.p).norm(),
		    nl = n.dot(r.d) < 0 ? n : n * -1, f = obj.c;
		float p = f.x > f.y && f.x > f.z ? f.x : f.y > f.z ? f.y : f.z; // max refl
		cl = cl + cf.mult(obj.e);
		if (++depth > 5) {
			if (rng() < p) {
				f = f * (1 / p);
			} else {
				return cl;
			}
		} // R.R.
		cf = cf.mult(f);
		if (obj.refl == DIFF) { // Ideal DIFFUSE reflection
			float r1 = 2 * M_PI * rng(), r2 = rng(), r2s = sqrt(r2);
			Vec w = nl,
			    u = ((fabs(w.x) > .1 ? Vec(0, 1) : Vec(1)) % w).norm(),
			    v = w % u;
			Vec d = (u * cos(r1) * r2s + v * sin(r1) * r2s +
					w * sqrt(1 - r2)).norm();
			r = Ray(x, d);
			continue;
		} else if (obj.refl == SPEC) { // Ideal SPECULAR reflection
			r = Ray(x, r.d - n * 2 * n.dot(r.d));
			continue;
		}
		Ray reflRay(x, r.d - n * 2 * n.dot(r.d)); // Ideal dielectric REFRACTION
		bool into = n.dot(nl) > 0;                // Ray from outside going in?
		float nc = 1, nt = 1.5, nnt = into ? nc / nt : nt / nc, ddn = r.d.dot(nl),
		      cos2t;
		if ((cos2t = 1 - nnt * nnt * (1 - ddn * ddn)) <
				0) { // Total internal reflection
			r = reflRay;
			continue;
		}
		Vec tdir =
			(r.d * nnt -
			 n * ((into ? 1 : -1) * (ddn * nnt + sqrt(cos2t)))).norm();
		float a = nt - nc, b = nt + nc, R0 = a * a / (b * b),
		      c = 1 - (into ? -ddn : tdir.dot(n));
		float Re = R0 + (1 - R0) * c * c * c * c * c, Tr = 1 - Re,
		      P = .25 + .5 * Re, RP = Re / P, TP = Tr / (1 - P);
		if (rng() < P) {
			cf = cf * RP;
			r = reflRay;
		} else {
			cf = cf * TP;
			r = Ray(x, tdir);
		}
		continue;
	}
}

__global__ void raytrace(Vec* c, const Sphere* spheres_, const int w, const int h, const int samps) {
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;
	Vec r;
	const Sphere *spheres = &spheres_[0];
	Ray cam(Vec(50, 52, 295.6), Vec(0, -0.042612, -1).norm()); // cam pos, dir
	Vec cx = Vec(w * .5135 / h), cy = (cx % cam.d).norm() * .5135;
	RNG rng(1 + (y * w) + x); // initialise our own rng with rand() seed
	for (int sy = 0, i = (h - y - 1) * w + x; sy < 2; sy++) // 2x2 subpixel rows
		for (int sx = 0; sx < 2; sx++, r = Vec()) {           // 2x2 subpixel cols
			for (int s = 0; s < samps; s++) {
				float r1 = 2 * rng(), dx = r1 < 1 ? sqrt(r1) - 1
					: 1 - sqrt(2 - r1);
				float r2 = 2 * rng(), dy = r2 < 1 ? sqrt(r2) - 1
					: 1 - sqrt(2 - r2);
				Vec d = cx * (((sx + .5 + dx) / 2 + x) / w - .5) +
					cy * (((sy + .5 + dy) / 2 + y) / h - .5) + cam.d;
				r = r + radiance(Ray(cam.o + d * 140, d.norm()), 0, spheres, rng) *
					(1. / samps);
			} // Camera rays are pushed ^^^^^ forward to start in interior
			c[i] = c[i] + Vec(clamp(r.x), clamp(r.y), clamp(r.z)) * .25;
		}
};

int main(int argc, char *argv[]) {

	const Sphere spheres_glob[] = {
		// Scene: radius, position, emission, color, material
		Sphere(1e4, Vec(1e4 + 1, 40.8, 81.6), Vec(), Vec(.75, .25, .25),
				DIFF), // Left
		Sphere(1e4, Vec(-1e4 + 99, 40.8, 81.6), Vec(), Vec(.25, .25, .75),
				DIFF),                                                     // Rght
		Sphere(1e4, Vec(50, 40.8, 1e4), Vec(), Vec(.75, .75, .75), DIFF), // Back
		Sphere(1e4, Vec(50, 40.8, -1e4 + 170), Vec(), Vec(), DIFF),       // Frnt
		Sphere(1e4, Vec(50, 1e4, 81.6), Vec(), Vec(.75, .75, .75), DIFF), // Botm
		Sphere(1e4, Vec(50, -1e4 + 81.6, 81.6), Vec(), Vec(.75, .75, .75),
				DIFF),                                                      // Top
		Sphere(16.5, Vec(27, 16.5, 47), Vec(), Vec(1, 1, 1) * .999, SPEC), // Mirr
		Sphere(16.5, Vec(73, 16.5, 78), Vec(), Vec(1, 1, 1) * .999, REFR), // Glas
		Sphere(600, Vec(50, 681.6 - .27, 81.6), Vec(12, 12, 12), Vec(),
				DIFF) // Lite
	};

	int w = 1024, h = 768; 
	int samps = 40; // # samples
	Vec *c = (Vec*) malloc(sizeof(Vec) * w * h);

	Vec *color_buffer;
	Sphere *spheres_buffer;

	hipMalloc((void**)&color_buffer, sizeof(Vec) * w * h);
	hipMalloc((void**)&spheres_buffer, sizeof(Sphere) * 9);

	hipMemcpy(spheres_buffer, &spheres_glob[0], sizeof(Sphere) * 9, hipMemcpyHostToDevice);

	raytrace<<< dim3(w/8, h/8), dim3(8, 8) >>>(color_buffer, spheres_buffer, w, h, samps);

	hipMemcpy(c, color_buffer, sizeof(Vec) * w * h, hipMemcpyDeviceToHost);

	FILE *f = fopen("image.ppm", "w"); // Write image to PPM file.
	fprintf(f, "P3\n%d %d\n%d\n", w, h, 255);
	for (int i = 0; i < w * h; i++)
		fprintf(f, "%d %d %d ", toInt(c[i].x), toInt(c[i].y), toInt(c[i].z));

	free(c);
	hipFree(color_buffer);
	hipFree(spheres_buffer);
	return 0;
}

	.text
	.file	"main.hip"
	.globl	_Z23__device_stub__raytraceP3VecPK6Sphereiii # -- Begin function _Z23__device_stub__raytraceP3VecPK6Sphereiii
	.type	_Z23__device_stub__raytraceP3VecPK6Sphereiii,@function
_Z23__device_stub__raytraceP3VecPK6Sphereiii: # @_Z23__device_stub__raytraceP3VecPK6Sphereiii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z8raytraceP3VecPK6Sphereiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z23__device_stub__raytraceP3VecPK6Sphereiii, .Lfunc_end0-_Z23__device_stub__raytraceP3VecPK6Sphereiii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0x461c4400                      # float 10001
	.long	0x42233333                      # float 40.7999992
	.zero	4
	.zero	4
.LCPI1_1:
	.long	0x3f400000                      # float 0.75
	.long	0x3e800000                      # float 0.25
	.zero	4
	.zero	4
.LCPI1_2:
	.long	0xc61ab400                      # float -9901
	.long	0x42233333                      # float 40.7999992
	.zero	4
	.zero	4
.LCPI1_3:
	.long	0x3e800000                      # float 0.25
	.long	0x3e800000                      # float 0.25
	.zero	4
	.zero	4
.LCPI1_4:
	.long	0x42480000                      # float 50
	.long	0x42233333                      # float 40.7999992
	.zero	4
	.zero	4
.LCPI1_5:
	.long	0x3f400000                      # float 0.75
	.long	0x3f400000                      # float 0.75
	.zero	4
	.zero	4
.LCPI1_6:
	.long	0x42480000                      # float 50
	.long	0x461c4000                      # float 1.0E+4
	.zero	4
	.zero	4
.LCPI1_7:
	.long	0x42480000                      # float 50
	.long	0xc61af99a                      # float -9918.40039
	.zero	4
	.zero	4
.LCPI1_8:
	.long	0x41d80000                      # float 27
	.long	0x41840000                      # float 16.5
	.zero	4
	.zero	4
.LCPI1_9:
	.long	0x3f7fbe77                      # float 0.999000012
	.long	0x3f7fbe77                      # float 0.999000012
	.zero	4
	.zero	4
.LCPI1_10:
	.long	0x42920000                      # float 73
	.long	0x41840000                      # float 16.5
	.zero	4
	.zero	4
.LCPI1_11:
	.long	0x42480000                      # float 50
	.long	0x442a551f                      # float 681.330017
	.zero	4
	.zero	4
.LCPI1_12:
	.long	0x41400000                      # float 12
	.long	0x41400000                      # float 12
	.zero	4
	.zero	4
.LCPI1_17:
	.zero	16
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI1_13:
	.long	0x3f800000                      # float 1
.LCPI1_18:
	.long	0x00000000                      # float 0
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI1_14:
	.quad	0x3fdd1745c0000000              # double 0.45454543828964233
.LCPI1_15:
	.quad	0x406fe00000000000              # double 255
.LCPI1_16:
	.quad	0x3fe0000000000000              # double 0.5
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$416, %rsp                      # imm = 0x1A0
	.cfi_def_cfa_offset 464
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$1176256512, %edx               # imm = 0x461C4000
	leaq	16(%rsp), %r14
	movl	%edx, (%r14)
	movsd	.LCPI1_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 4(%r14)
	movl	$1117991731, %eax               # imm = 0x42A33333
	movl	%eax, 12(%r14)
	xorl	%ecx, %ecx
	movq	%rcx, 16(%r14)
	movl	%ecx, 24(%r14)
	movsd	.LCPI1_1(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 28(%r14)
	movq	$1048576000, 36(%r14)           # imm = 0x3E800000
	movl	%edx, 44(%r14)
	movsd	.LCPI1_2(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 48(%r14)
	movl	%eax, 56(%r14)
	movq	%rcx, 60(%r14)
	movl	%ecx, 68(%r14)
	movsd	.LCPI1_3(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 72(%r14)
	movl	$1061158912, %esi               # imm = 0x3F400000
	movq	%rsi, 80(%r14)
	movl	%edx, 88(%r14)
	movsd	.LCPI1_4(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 92(%r14)
	movl	%edx, 100(%r14)
	movq	%rcx, 104(%r14)
	movl	%ecx, 112(%r14)
	movsd	.LCPI1_5(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	%xmm1, 116(%r14)
	movq	%rsi, 124(%r14)
	movl	%edx, 132(%r14)
	movsd	%xmm0, 136(%r14)
	movl	$-971401216, 144(%r14)          # imm = 0xC6199800
	xorps	%xmm0, %xmm0
	movups	%xmm0, 160(%r14)
	movups	%xmm0, 148(%r14)
	movl	%edx, 176(%r14)
	movsd	.LCPI1_6(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 180(%r14)
	movl	%eax, 188(%r14)
	movq	%rcx, 192(%r14)
	movl	%ecx, 200(%r14)
	movsd	%xmm1, 204(%r14)
	movq	%rsi, 212(%r14)
	movl	%edx, 220(%r14)
	movsd	.LCPI1_7(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 224(%r14)
	movl	%eax, 232(%r14)
	movq	%rcx, 236(%r14)
	movl	%ecx, 244(%r14)
	movsd	%xmm1, 248(%r14)
	movq	%rsi, 256(%r14)
	movl	$1099169792, %edx               # imm = 0x41840000
	movl	%edx, 264(%r14)
	movsd	.LCPI1_8(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 268(%r14)
	movl	$1111228416, 276(%r14)          # imm = 0x423C0000
	movq	%rcx, 280(%r14)
	movl	%ecx, 288(%r14)
	movsd	.LCPI1_9(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	%xmm0, 292(%r14)
	movabsq	$5360303735, %rsi               # imm = 0x13F7FBE77
	movq	%rsi, 300(%r14)
	movl	%edx, 308(%r14)
	movsd	.LCPI1_10(%rip), %xmm1          # xmm1 = mem[0],zero
	movsd	%xmm1, 312(%r14)
	movl	$1117519872, 320(%r14)          # imm = 0x429C0000
	movq	%rcx, 324(%r14)
	movl	%ecx, 332(%r14)
	movsd	%xmm0, 336(%r14)
	movabsq	$9655271031, %rcx               # imm = 0x23F7FBE77
	movq	%rcx, 344(%r14)
	movl	$1142292480, 352(%r14)          # imm = 0x44160000
	movsd	.LCPI1_11(%rip), %xmm0          # xmm0 = mem[0],zero
	movsd	%xmm0, 356(%r14)
	movl	%eax, 364(%r14)
	movsd	.LCPI1_12(%rip), %xmm0          # xmm0 = mem[0],zero
	movsd	%xmm0, 368(%r14)
	movl	$1094713344, 376(%r14)          # imm = 0x41400000
	movl	$9437184, %edi                  # imm = 0x900000
	callq	malloc
	movq	%rax, %rbx
	xorpd	%xmm0, %xmm0
	movupd	%xmm0, 380(%r14)
	movq	%rsp, %rdi
	movl	$9437184, %esi                  # imm = 0x900000
	callq	hipMalloc
	leaq	8(%rsp), %r15
	movl	$396, %esi                      # imm = 0x18C
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$396, %edx                      # imm = 0x18C
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$412316860544, %rdi             # imm = 0x6000000080
	movabsq	$34359738376, %rdx              # imm = 0x800000008
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$1024, %edx                     # imm = 0x400
	movl	$768, %ecx                      # imm = 0x300
	movl	$40, %r8d
	callq	_Z23__device_stub__raytraceP3VecPK6Sphereiii
.LBB1_2:
	movq	(%rsp), %rsi
	movl	$9437184, %edx                  # imm = 0x900000
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %r14
	movl	$.L.str.2, %esi
	movq	%rax, %rdi
	movl	$1024, %edx                     # imm = 0x400
	movl	$768, %ecx                      # imm = 0x300
	movl	$255, %r8d
	xorl	%eax, %eax
	callq	fprintf
	movl	$8, %r12d
.LBB1_3:                                # =>This Inner Loop Header: Depth=1
	movss	-8(%rbx,%r12), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI1_13(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	minss	%xmm0, %xmm1
	xorps	%xmm2, %xmm2
	cmpltss	%xmm2, %xmm0
	andnps	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	.LCPI1_14(%rip), %xmm1          # xmm1 = mem[0],zero
	callq	pow
	movsd	.LCPI1_15(%rip), %xmm1          # xmm1 = mem[0],zero
	mulsd	%xmm1, %xmm0
	movsd	.LCPI1_16(%rip), %xmm1          # xmm1 = mem[0],zero
	addsd	%xmm1, %xmm0
	cvttsd2si	%xmm0, %ebp
	movss	-4(%rbx,%r12), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI1_13(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	minss	%xmm0, %xmm1
	cmpltss	.LCPI1_18(%rip), %xmm0
	andnps	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	.LCPI1_14(%rip), %xmm1          # xmm1 = mem[0],zero
	callq	pow
	mulsd	.LCPI1_15(%rip), %xmm0
	addsd	.LCPI1_16(%rip), %xmm0
	cvttsd2si	%xmm0, %r15d
	movss	(%rbx,%r12), %xmm0              # xmm0 = mem[0],zero,zero,zero
	movss	.LCPI1_13(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	minss	%xmm0, %xmm1
	cmpltss	.LCPI1_18(%rip), %xmm0
	andnps	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	.LCPI1_14(%rip), %xmm1          # xmm1 = mem[0],zero
	callq	pow
	mulsd	.LCPI1_15(%rip), %xmm0
	addsd	.LCPI1_16(%rip), %xmm0
	cvttsd2si	%xmm0, %r8d
	movl	$.L.str.3, %esi
	movq	%r14, %rdi
	movl	%ebp, %edx
	movl	%r15d, %ecx
	xorl	%eax, %eax
	callq	fprintf
	addq	$12, %r12
	cmpq	$9437192, %r12                  # imm = 0x900008
	jne	.LBB1_3
# %bb.4:
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$416, %rsp                      # imm = 0x1A0
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8raytraceP3VecPK6Sphereiii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z8raytraceP3VecPK6Sphereiii,@object # @_Z8raytraceP3VecPK6Sphereiii
	.section	.rodata,"a",@progbits
	.globl	_Z8raytraceP3VecPK6Sphereiii
	.p2align	3, 0x0
_Z8raytraceP3VecPK6Sphereiii:
	.quad	_Z23__device_stub__raytraceP3VecPK6Sphereiii
	.size	_Z8raytraceP3VecPK6Sphereiii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"image.ppm"
	.size	.L.str, 10

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"w"
	.size	.L.str.1, 2

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"P3\n%d %d\n%d\n"
	.size	.L.str.2, 13

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"%d %d %d "
	.size	.L.str.3, 10

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z8raytraceP3VecPK6Sphereiii"
	.size	.L__unnamed_1, 29

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.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 _Z23__device_stub__raytraceP3VecPK6Sphereiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z8raytraceP3VecPK6Sphereiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8raytraceP3VecPK6Sphereiii ; -- Begin function _Z8raytraceP3VecPK6Sphereiii
	.globl	_Z8raytraceP3VecPK6Sphereiii
	.p2align	8
	.type	_Z8raytraceP3VecPK6Sphereiii,@function
_Z8raytraceP3VecPK6Sphereiii:           ; @_Z8raytraceP3VecPK6Sphereiii
; %bb.0:
	s_load_b256 s[4:11], s[0:1], 0x0
	s_mov_b32 s2, 0x8d4fdf3b
	s_mov_b32 s3, 0x3fe06e97
	s_load_b32 s0, s[0:1], 0x2c
	v_mov_b32_e32 v32, 0x3ee38e3a
	s_mov_b32 s12, 0x54442d18
	s_mov_b32 s18, 0x430c0000
	s_mov_b32 s13, 0x401921fb
	s_mov_b32 s19, 0x800000
	s_mov_b32 s20, 0xb94c1982
	s_mov_b32 s21, 0x37d75334
	v_mov_b32_e32 v31, 0x4393cccd
	v_mov_b32_e32 v33, 0x3f2aaaab
	s_waitcnt lgkmcnt(0)
	v_cvt_f64_i32_e32 v[11:12], s8
	v_cvt_f64_i32_e32 v[13:14], s9
	v_cvt_f64_i32_e32 v[19:20], s10
	s_lshr_b32 s1, s0, 16
	s_and_b32 s0, s0, 0xffff
	s_cmp_gt_i32 s10, 0
	s_cselect_b32 s11, -1, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[1:2], v[11:12], s[2:3]
	v_div_scale_f64 v[21:22], null, v[19:20], v[19:20], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_scale_f64 v[3:4], null, v[13:14], v[13:14], v[1:2]
	v_div_scale_f64 v[9:10], vcc_lo, v[1:2], v[13:14], v[1:2]
	v_rcp_f64_e32 v[17:18], v[21:22]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[5:6], v[3:4]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[7:8], -v[3:4], v[5:6], 1.0
	v_fma_f64 v[5:6], v[5:6], v[7:8], v[5:6]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[7:8], -v[3:4], v[5:6], 1.0
	v_fma_f64 v[5:6], v[5:6], v[7:8], v[5:6]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[7:8], v[9:10], v[5:6]
	v_fma_f64 v[3:4], -v[3:4], v[7:8], v[9:10]
	v_bfe_u32 v9, v0, 10, 10
	v_and_b32_e32 v0, 0x3ff, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[23:24], null, s15, s1, v[9:10]
	v_div_fmas_f64 v[3:4], v[3:4], v[5:6], v[7:8]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[5:6], null, s14, s0, v[0:1]
	v_fma_f64 v[6:7], -v[21:22], v[17:18], 1.0
	v_xad_u32 v0, v23, -1, s9
	s_mov_b32 s14, 0x9999999a
	s_mov_b32 s15, 0x3fb99999
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mad_u64_u32 v[8:9], null, v0, s8, v[5:6]
	v_div_fixup_f64 v[0:1], v[3:4], v[13:14], v[1:2]
	v_fma_f64 v[2:3], v[17:18], v[6:7], v[17:18]
	v_div_scale_f64 v[6:7], s0, 1.0, v[19:20], 1.0
	s_delay_alu instid0(VALU_DEP_4)
	v_mad_i64_i32 v[15:16], null, v8, 12, s[4:5]
	s_add_u32 s4, s6, 0x16c
	s_addc_u32 s5, s7, 0
	global_load_b96 v[8:10], v[15:16], off
	v_cvt_f32_f64_e32 v25, v[0:1]
	v_fma_f64 v[0:1], -v[21:22], v[2:3], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v18, 0 :: v_dual_mul_f32 v17, 0xbd2e6169, v25
	v_fmamk_f32 v4, v25, 0x3f7fc494, v18
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[0:1], v[2:3], v[0:1], v[2:3]
	v_mul_f32_e32 v2, v4, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v2, v17, v17
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v2
	v_mul_f32_e32 v3, 0x4f800000, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v24, v2, v3, vcc_lo
	v_sqrt_f32_e32 v26, v24
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v27, -1, v26
	v_add_nc_u32_e32 v28, 1, v26
	v_mul_f64 v[2:3], v[6:7], v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v29, -v27, v26, v24
	v_fma_f32 v30, -v28, v26, v24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s1, 0, v29
	v_cndmask_b32_e64 v26, v26, v27, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cmp_lt_f32_e64 s1, 0, v30
	v_mov_b32_e32 v30, 0x42500000
	v_cndmask_b32_e64 v26, v26, v28, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v27, 0x37800000, v26
	v_cndmask_b32_e32 v26, v26, v27, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v24, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v24, v26, v24, vcc_lo
	v_fma_f64 v[6:7], -v[21:22], v[2:3], v[6:7]
	v_div_scale_f32 v26, null, v24, v24, 1.0
	v_div_scale_f32 v22, vcc_lo, 1.0, v24, 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v27, v26
	s_waitcnt_depctr 0xfff
	v_fma_f32 v21, -v26, v27, 1.0
	v_fmac_f32_e32 v27, v21, v27
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v21, v22, v27
	v_fma_f32 v28, -v26, v21, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v21, v28, v27
	v_fma_f32 v22, -v26, v21, v22
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fmas_f32 v27, v22, v27, v21
	s_mov_b32 vcc_lo, s0
	v_cvt_f64_i32_e32 v[21:22], v23
	v_div_fmas_f64 v[0:1], v[6:7], v[0:1], v[2:3]
	v_mul_lo_u32 v3, v23, s8
	s_mov_b32 s0, -1
	s_mov_b64 s[8:9], 0.5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[0:1], v[0:1], v[19:20], 1.0
	v_cvt_f64_i32_e32 v[19:20], v5
	v_cvt_f32_f64_e32 v26, v[0:1]
	v_div_fixup_f32 v0, v27, v24, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v1, 0, v0
	v_mul_f32_e32 v2, v4, v0
	v_dual_mul_f32 v0, v17, v0 :: v_dual_mul_f32 v27, 0x3f0374bc, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mul_f32 v28, 0x3f0374bc, v2 :: v_dual_mul_f32 v29, 0x3f0374bc, v0
	v_add3_u32 v0, v5, v3, 1
.LBB0_1:                                ; %.preheader170
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_2 Depth 2
                                        ;       Child Loop BB0_4 Depth 3
                                        ;         Child Loop BB0_13 Depth 4
                                        ;           Child Loop BB0_14 Depth 5
	s_xor_b32 s22, s0, -1
	s_mov_b32 s23, -1
	s_mov_b64 s[16:17], 0.5
.LBB0_2:                                ; %.preheader
                                        ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_4 Depth 3
                                        ;         Child Loop BB0_13 Depth 4
                                        ;           Child Loop BB0_14 Depth 5
	v_mov_b32_e32 v34, 0
	s_and_not1_b32 vcc_lo, exec_lo, s11
	s_cbranch_vccnz .LBB0_51
; %bb.3:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB0_2 Depth=2
	v_dual_mov_b32 v34, 0 :: v_dual_mov_b32 v35, 0
	v_mov_b32_e32 v36, 0
	s_mov_b32 s24, 0
.LBB0_4:                                ; %.lr.ph
                                        ;   Parent Loop BB0_1 Depth=1
                                        ;     Parent Loop BB0_2 Depth=2
                                        ; =>    This Loop Header: Depth=3
                                        ;         Child Loop BB0_13 Depth 4
                                        ;           Child Loop BB0_14 Depth 5
	v_lshrrev_b32_e32 v1, 6, v0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v1, v0
	v_lshlrev_b32_e32 v1, 17, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v1, v0
	v_lshrrev_b32_e32 v1, 9, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v1, v0
	v_and_or_b32 v0, 0x7fffff, v1, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v0, -1.0, v0
	v_add_f32_e32 v2, v0, v0
                                        ; implicit-def: $vgpr0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ngt_f32_e32 1.0, v2
	s_xor_b32 s1, exec_lo, s0
	s_cbranch_execz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_4 Depth=3
	v_sub_f32_e32 v0, 2.0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v2, 0x4f800000, v0
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v0
	v_cndmask_b32_e32 v0, v0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v2, v0
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v3, -1, v2
	v_add_nc_u32_e32 v4, 1, v2
	v_fma_f32 v5, -v3, v2, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v2, v0
	v_cmp_ge_f32_e64 s0, 0, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v2, v2, v3, s0
	v_cmp_lt_f32_e64 s0, 0, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v2, v2, v4, s0
	v_mul_f32_e32 v3, 0x37800000, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v0, 0x260
	v_cndmask_b32_e32 v0, v2, v0, vcc_lo
                                        ; implicit-def: $vgpr2
	s_delay_alu instid0(VALU_DEP_1)
	v_sub_f32_e32 v0, 1.0, v0
.LBB0_6:                                ; %Flow269
                                        ;   in Loop: Header=BB0_4 Depth=3
	s_and_not1_saveexec_b32 s1, s1
	s_cbranch_execz .LBB0_8
; %bb.7:                                ;   in Loop: Header=BB0_4 Depth=3
	v_mul_f32_e32 v0, 0x4f800000, v2
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v0, v2, v0, vcc_lo
	v_sqrt_f32_e32 v2, v0
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v3, -1, v2
	v_add_nc_u32_e32 v4, 1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v5, -v3, v2, v0
	v_fma_f32 v6, -v4, v2, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v5
	v_cndmask_b32_e64 v2, v2, v3, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v6
	v_cndmask_b32_e64 v2, v2, v4, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v3, 0x37800000, v2
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v0, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v0, v2, v0, vcc_lo
	v_add_f32_e32 v0, -1.0, v0
.LBB0_8:                                ;   in Loop: Header=BB0_4 Depth=3
	s_or_b32 exec_lo, exec_lo, s1
	v_lshrrev_b32_e32 v2, 6, v1
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v2, v1
	v_lshlrev_b32_e32 v2, 17, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v2, v1
	v_lshrrev_b32_e32 v2, 9, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v45, v2, v1
	v_and_or_b32 v1, 0x7fffff, v45, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v1, -1.0, v1
	v_add_f32_e32 v2, v1, v1
                                        ; implicit-def: $vgpr1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ngt_f32_e32 1.0, v2
	s_xor_b32 s1, exec_lo, s0
	s_cbranch_execz .LBB0_10
; %bb.9:                                ;   in Loop: Header=BB0_4 Depth=3
	v_sub_f32_e32 v1, 2.0, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v2, 0x4f800000, v1
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v1
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v2, v1
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v3, -1, v2
	v_add_nc_u32_e32 v4, 1, v2
	v_fma_f32 v5, -v3, v2, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v2, v1
	v_cmp_ge_f32_e64 s0, 0, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v2, v2, v3, s0
	v_cmp_lt_f32_e64 s0, 0, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v2, v2, v4, s0
	v_mul_f32_e32 v3, 0x37800000, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v1, 0x260
	v_cndmask_b32_e32 v1, v2, v1, vcc_lo
                                        ; implicit-def: $vgpr2
	s_delay_alu instid0(VALU_DEP_1)
	v_sub_f32_e32 v1, 1.0, v1
.LBB0_10:                               ; %Flow268
                                        ;   in Loop: Header=BB0_4 Depth=3
	s_and_not1_saveexec_b32 s1, s1
	s_cbranch_execz .LBB0_12
; %bb.11:                               ;   in Loop: Header=BB0_4 Depth=3
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v2
	v_mul_f32_e32 v1, 0x4f800000, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v2, v1, vcc_lo
	v_sqrt_f32_e32 v2, v1
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v4, 1, v2
	v_add_nc_u32_e32 v3, -1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v2, v1
	v_fma_f32 v5, -v3, v2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v5
	v_cndmask_b32_e64 v2, v2, v3, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v6
	v_cndmask_b32_e64 v2, v2, v4, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v3, 0x37800000, v2
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v1, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v2, v1, vcc_lo
	v_add_f32_e32 v1, -1.0, v1
.LBB0_12:                               ;   in Loop: Header=BB0_4 Depth=3
	s_or_b32 exec_lo, exec_lo, s1
	v_cvt_f64_f32_e32 v[2:3], v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cvt_f64_f32_e32 v[0:1], v1
	s_mov_b32 s25, 0
	v_mov_b32_e32 v48, 1.0
	v_add_f64 v[2:3], s[16:17], v[2:3]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[0:1], s[8:9], v[0:1]
	v_fma_f64 v[2:3], v[2:3], 0.5, v[19:20]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[0:1], v[0:1], 0.5, v[21:22]
	v_div_scale_f64 v[4:5], null, v[11:12], v[11:12], v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_scale_f64 v[6:7], null, v[13:14], v[13:14], v[0:1]
	v_div_scale_f64 v[43:44], vcc_lo, v[2:3], v[11:12], v[2:3]
	v_rcp_f64_e32 v[23:24], v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[37:38], v[6:7]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[39:40], -v[4:5], v[23:24], 1.0
	v_fma_f64 v[41:42], -v[6:7], v[37:38], 1.0
	v_fma_f64 v[23:24], v[23:24], v[39:40], v[23:24]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[37:38], v[37:38], v[41:42], v[37:38]
	v_fma_f64 v[39:40], -v[4:5], v[23:24], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[41:42], -v[6:7], v[37:38], 1.0
	v_fma_f64 v[23:24], v[23:24], v[39:40], v[23:24]
	v_div_scale_f64 v[39:40], s0, v[0:1], v[13:14], v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[37:38], v[37:38], v[41:42], v[37:38]
	v_mul_f64 v[41:42], v[43:44], v[23:24]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[46:47], v[39:40], v[37:38]
	v_fma_f64 v[4:5], -v[4:5], v[41:42], v[43:44]
	v_mov_b32_e32 v44, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f64 v[6:7], -v[6:7], v[46:47], v[39:40]
	v_dual_mov_b32 v39, 0 :: v_dual_mov_b32 v40, 0
	v_div_fmas_f64 v[4:5], v[4:5], v[23:24], v[41:42]
	s_mov_b32 vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_div_fmas_f64 v[6:7], v[6:7], v[37:38], v[46:47]
	v_dual_mov_b32 v46, 1.0 :: v_dual_mov_b32 v47, 1.0
	v_dual_mov_b32 v37, 0 :: v_dual_mov_b32 v38, 0
	v_div_fixup_f64 v[2:3], v[4:5], v[11:12], v[2:3]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[0:1], v[6:7], v[13:14], v[0:1]
	v_add_f64 v[2:3], v[2:3], -0.5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[0:1], v[0:1], -0.5
	v_cvt_f32_f64_e32 v2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f32_f64_e32 v0, v[0:1]
	v_mul_f32_e32 v1, 0, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v3, v27, v0
	v_fma_f32 v4, v28, v0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v3, v25, v2 :: v_dual_add_f32 v2, 0xbd2e6169, v4
	v_dual_fmac_f32 v1, v29, v0 :: v_dual_add_f32 v0, 0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v3, v2, v2
	v_add_f32_e32 v1, 0xbf7fc494, v1
	v_fmamk_f32 v41, v2, 0x430c0000, v30
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmaak_f32 v43, s18, v0, 0x42480000
	v_fmac_f32_e32 v3, v0, v0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmamk_f32 v42, v1, 0x430c0000, v31
	v_fmac_f32_e32 v3, v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v4, 0x4f800000, v3
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v3
	v_cndmask_b32_e32 v3, v3, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v5, -1, v4
	v_add_nc_u32_e32 v6, 1, v4
	v_fma_f32 v7, -v5, v4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v17, -v6, v4, v3
	v_cmp_ge_f32_e64 s0, 0, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v4, v4, v5, s0
	v_cmp_lt_f32_e64 s0, 0, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v4, v4, v6, s0
	v_mul_f32_e32 v5, 0x37800000, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v3, 0x260
	v_cndmask_b32_e32 v3, v4, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v4, null, v3, v3, 1.0
	v_div_scale_f32 v7, vcc_lo, 1.0, v3, 1.0
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v4, v5, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v6, v5
	v_mul_f32_e32 v6, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v17, -v4, v6, v7
	v_fmac_f32_e32 v6, v17, v5
	v_mov_b32_e32 v17, v45
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v4, -v4, v6, v7
	v_div_fmas_f32 v4, v4, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v3, v4, v3, 1.0
	v_mul_f32_e32 v50, v0, v3
	v_mul_f32_e32 v51, v2, v3
	v_mul_f32_e32 v49, v1, v3
.LBB0_13:                               ; %.critedge.i
                                        ;   Parent Loop BB0_1 Depth=1
                                        ;     Parent Loop BB0_2 Depth=2
                                        ;       Parent Loop BB0_4 Depth=3
                                        ; =>      This Loop Header: Depth=4
                                        ;           Child Loop BB0_14 Depth 5
	v_mov_b32_e32 v52, 0x60ad78ec
	s_mov_b32 s1, 8
	s_mov_b64 s[2:3], s[4:5]
.LBB0_14:                               ;   Parent Loop BB0_1 Depth=1
                                        ;     Parent Loop BB0_2 Depth=2
                                        ;       Parent Loop BB0_4 Depth=3
                                        ;         Parent Loop BB0_13 Depth=4
                                        ; =>        This Inner Loop Header: Depth=5
	global_load_b128 v[1:4], v18, s[2:3] offset:-12
	s_mov_b32 s26, exec_lo
	s_waitcnt vmcnt(0)
	v_dual_sub_f32 v3, v3, v41 :: v_dual_sub_f32 v2, v2, v43
	v_sub_f32_e32 v4, v4, v42
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v0, v51, v3
	v_mul_f32_e32 v3, v3, v3
	v_fmac_f32_e32 v3, v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v0, v50, v2 :: v_dual_fmac_f32 v3, v4, v4
	v_fmac_f32_e32 v0, v49, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v2, v0, v0, -v3
	v_dual_fmac_f32 v2, v1, v1 :: v_dual_mov_b32 v1, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ngt_f32_e32 0, v2
	s_cbranch_execz .LBB0_16
; %bb.15:                               ;   in Loop: Header=BB0_14 Depth=5
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v2
	v_mul_f32_e32 v1, 0x4f800000, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v2, v1, vcc_lo
	v_sqrt_f32_e32 v2, v1
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v4, 1, v2
	v_add_nc_u32_e32 v3, -1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v2, v1
	v_fma_f32 v5, -v3, v2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v5
	v_cndmask_b32_e64 v2, v2, v3, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v6
	v_cndmask_b32_e64 v2, v2, v4, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v3, 0x37800000, v2
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v1, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v2, v1, vcc_lo
	v_add_f32_e32 v2, v0, v1
	v_sub_f32_e32 v0, v0, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cmp_lt_f32_e32 vcc_lo, 0x3c75c28f, v2
	v_cndmask_b32_e32 v1, 0, v2, vcc_lo
	v_cmp_lt_f32_e32 vcc_lo, 0x3c75c28f, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_cndmask_b32_e32 v1, v1, v0, vcc_lo
.LBB0_16:                               ; %_ZNK6Sphere9intersectERK3Ray.exit.i.i
                                        ;   in Loop: Header=BB0_14 Depth=5
	s_or_b32 exec_lo, exec_lo, s26
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_neq_f32_e32 vcc_lo, 0, v1
	v_cmp_lt_f32_e64 s0, v1, v52
	s_and_b32 vcc_lo, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cndmask_b32_e64 v37, v37, s1, vcc_lo
	s_add_i32 s1, s1, -1
	v_cndmask_b32_e32 v52, v52, v1, vcc_lo
	s_add_u32 s2, s2, 0xffffffd4
	s_addc_u32 s3, s3, -1
	s_cmp_lg_u32 s1, -1
	s_cbranch_scc1 .LBB0_14
; %bb.17:                               ; %_Z9intersectIPK6SphereEbRK3RayRfRiT_.exit.i
                                        ;   in Loop: Header=BB0_13 Depth=4
	v_mov_b32_e32 v0, v45
	s_mov_b32 s0, -1
	s_mov_b32 s3, exec_lo
	v_cmpx_gt_f32_e32 0x60ad78ec, v52
	s_cbranch_execz .LBB0_49
; %bb.18:                               ;   in Loop: Header=BB0_13 Depth=4
	v_mad_i64_i32 v[23:24], null, v37, 44, s[6:7]
	s_mov_b32 s2, -1
	s_mov_b32 s26, exec_lo
	s_clause 0x2
	global_load_b128 v[4:7], v[23:24], off offset:4
	global_load_b128 v[0:3], v[23:24], off offset:20
	global_load_b32 v56, v[23:24], off offset:36
	s_waitcnt vmcnt(1)
	v_dual_fmac_f32 v40, v47, v7 :: v_dual_fmac_f32 v39, v46, v1
	v_fmac_f32_e32 v38, v48, v0
	v_cmpx_lt_u32_e32 4, v44
	s_cbranch_execz .LBB0_22
; %bb.19:                               ;   in Loop: Header=BB0_13 Depth=4
	v_lshrrev_b32_e32 v0, 6, v17
	v_cmp_gt_f32_e32 vcc_lo, v2, v3
	s_waitcnt vmcnt(0)
	v_cmp_gt_f32_e64 s0, v2, v56
	v_cmp_gt_f32_e64 s1, v3, v56
	v_xor_b32_e32 v0, v0, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_and_b32 vcc_lo, vcc_lo, s0
	s_mov_b32 s0, 0
	v_lshlrev_b32_e32 v1, 17, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v1, v0
	v_lshrrev_b32_e32 v1, 9, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v17, v1, v0
	v_cndmask_b32_e64 v0, v56, v3, s1
	s_mov_b32 s1, exec_lo
	v_and_or_b32 v1, 0x7fffff, v17, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v0, v0, v2 :: v_dual_add_f32 v1, -1.0, v1
	v_cmpx_lt_f32_e32 v1, v0
	s_cbranch_execz .LBB0_21
; %bb.20:                               ;   in Loop: Header=BB0_13 Depth=4
	v_div_scale_f32 v1, null, v0, v0, 1.0
	v_div_scale_f32 v53, vcc_lo, 1.0, v0, 1.0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v7, v1
	s_waitcnt_depctr 0xfff
	v_fma_f32 v45, -v1, v7, 1.0
	v_fmac_f32_e32 v7, v45, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v45, v53, v7
	v_fma_f32 v54, -v1, v45, v53
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v45, v54, v7
	v_fma_f32 v1, -v1, v45, v53
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v1, v1, v7, v45
	v_div_fixup_f32 v0, v1, v0, 1.0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f32_e32 v56, v56, v0
	v_mul_f32_e32 v2, v2, v0
	v_mul_f32_e32 v3, v3, v0
.LBB0_21:                               ; %Flow266
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v45, v17
	s_or_not1_b32 s0, s0, exec_lo
.LBB0_22:                               ; %Flow265
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s26
	s_and_saveexec_b32 s26, s0
	s_cbranch_execz .LBB0_48
; %bb.23:                               ; %NodeBlock
                                        ;   in Loop: Header=BB0_13 Depth=4
	global_load_b32 v57, v[23:24], off offset:40
	v_fmac_f32_e32 v43, v50, v52
	v_fmac_f32_e32 v42, v49, v52
	v_add_nc_u32_e32 v44, 1, v44
	s_mov_b32 s2, 0
	s_mov_b32 s1, 0
	v_sub_f32_e32 v1, v43, v4
	v_fmac_f32_e32 v41, v51, v52
	s_mov_b32 s27, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v0, v41, v5 :: v_dual_sub_f32 v5, v42, v6
	v_mul_f32_e32 v4, v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v1, v1
	v_fmac_f32_e32 v4, v5, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v6, 0x4f800000, v4
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v4
	v_cndmask_b32_e32 v4, v4, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v6, v4
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v7, -1, v6
	v_add_nc_u32_e32 v23, 1, v6
	v_fma_f32 v24, -v7, v6, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v52, -v23, v6, v4
	v_cmp_ge_f32_e64 s0, 0, v24
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v6, v6, v7, s0
	v_cmp_lt_f32_e64 s0, 0, v52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v6, v6, v23, s0
	v_mul_f32_e32 v7, 0x37800000, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v6, v6, v7, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v4, 0x260
	v_cndmask_b32_e32 v4, v6, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f32 v6, null, v4, v4, 1.0
	v_rcp_f32_e32 v7, v6
	s_waitcnt_depctr 0xfff
	v_fma_f32 v23, -v6, v7, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v7, v23, v7
	v_div_scale_f32 v23, vcc_lo, 1.0, v4, 1.0
	v_mul_f32_e32 v24, v23, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v52, -v6, v24, v23
	v_fmac_f32_e32 v24, v52, v7
                                        ; implicit-def: $vgpr52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v6, -v6, v24, v23
                                        ; implicit-def: $vgpr23
	v_div_fmas_f32 v6, v6, v7, v24
                                        ; implicit-def: $vgpr24
                                        ; implicit-def: $vgpr7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v4, v6, v4, 1.0
	v_dual_mul_f32 v6, v1, v4 :: v_dual_mul_f32 v1, v48, v3
	v_mul_f32_e32 v54, v0, v4
	v_dual_mul_f32 v53, v5, v4 :: v_dual_mul_f32 v0, v47, v2
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v2, v46, v56
                                        ; implicit-def: $vgpr46
                                        ; implicit-def: $vgpr48
                                        ; implicit-def: $vgpr47
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v55, v51, v54
	v_fmac_f32_e32 v55, v50, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v55, v49, v53
	v_cmp_gt_f32_e32 vcc_lo, 0, v55
	s_waitcnt vmcnt(0)
	v_cmpx_lt_i32_e32 0, v57
	s_xor_b32 s27, exec_lo, s27
	s_cbranch_execz .LBB0_27
; %bb.24:                               ; %LeafBlock253
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_mov_b32 s1, -1
	s_mov_b32 s28, exec_lo
                                        ; implicit-def: $vgpr24
                                        ; implicit-def: $vgpr52
                                        ; implicit-def: $vgpr23
	v_cmpx_eq_u32_e32 1, v57
; %bb.25:                               ;   in Loop: Header=BB0_13 Depth=4
	v_add_f32_e32 v3, v6, v6
	v_dual_add_f32 v4, v54, v54 :: v_dual_add_f32 v5, v53, v53
	s_xor_b32 s1, exec_lo, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v23, -v3, v55, v50
	v_fma_f32 v52, -v4, v55, v51
	s_delay_alu instid0(VALU_DEP_3)
	v_fma_f32 v24, -v5, v55, v49
; %bb.26:                               ; %Flow260
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s28
	v_dual_mov_b32 v46, v2 :: v_dual_mov_b32 v47, v0
	v_mov_b32_e32 v48, v1
	v_mov_b32_e32 v7, v17
	s_and_b32 s1, s1, exec_lo
                                        ; implicit-def: $vgpr57
.LBB0_27:                               ; %Flow259
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_and_not1_saveexec_b32 s27, s27
; %bb.28:                               ; %LeafBlock
                                        ;   in Loop: Header=BB0_13 Depth=4
	v_cmp_ne_u32_e64 s0, 0, v57
	s_and_not1_b32 s1, s1, exec_lo
	s_mov_b32 s2, exec_lo
                                        ; implicit-def: $vgpr24
                                        ; implicit-def: $vgpr52
                                        ; implicit-def: $vgpr23
                                        ; implicit-def: $vgpr46
                                        ; implicit-def: $vgpr48
                                        ; implicit-def: $vgpr47
                                        ; implicit-def: $vgpr7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, exec_lo
	s_or_b32 s1, s1, s0
; %bb.29:                               ; %Flow261
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s27
	v_cndmask_b32_e64 v4, -v6, v6, vcc_lo
	v_cndmask_b32_e64 v5, -v54, v54, vcc_lo
	v_cndmask_b32_e64 v3, -v53, v53, vcc_lo
	s_and_saveexec_b32 s0, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s27, exec_lo, s0
	s_cbranch_execz .LBB0_37
; %bb.30:                               ;   in Loop: Header=BB0_13 Depth=4
	v_dual_mul_f32 v56, v51, v5 :: v_dual_add_f32 v23, v6, v6
	v_add_f32_e32 v58, v53, v53
	v_dual_add_f32 v52, v54, v54 :: v_dual_mov_b32 v47, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v56, v50, v4
	v_mul_f32_e32 v7, v54, v5
	v_mov_b32_e32 v46, v2
	v_mov_b32_e32 v48, v1
	v_fma_f32 v23, -v23, v55, v50
	v_dual_fmac_f32 v56, v49, v3 :: v_dual_fmac_f32 v7, v6, v4
	v_fma_f32 v52, -v52, v55, v51
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v24, -v56, v56, 1.0
	v_fmac_f32_e32 v7, v53, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v7
	v_cndmask_b32_e64 v7, 0x40100000, v32, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f32 v57, -v7, v24, 1.0
	v_mov_b32_e32 v7, v17
	v_fma_f32 v24, -v58, v55, v49
	v_cmpx_ngt_f32_e32 0, v57
	s_xor_b32 s28, exec_lo, s1
	s_cbranch_execz .LBB0_36
; %bb.31:                               ;   in Loop: Header=BB0_13 Depth=4
	v_mul_f32_e32 v7, 0x4f800000, v57
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v57
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v57, v7, vcc_lo
	v_sqrt_f32_e32 v45, v7
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v46, -1, v45
	v_add_nc_u32_e32 v47, 1, v45
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v48, -v46, v45, v7
	v_fma_f32 v55, -v47, v45, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s1, 0, v48
	v_cndmask_b32_e64 v45, v45, v46, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s1, 0, v55
	v_cndmask_b32_e64 v45, v45, v47, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v46, 0x37800000, v45
	v_cndmask_b32_e32 v45, v45, v46, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v7, 0x260
	v_cndmask_b32_e64 v46, 0x3fc00000, v33, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v7, v45, v7, vcc_lo
	v_fmac_f32_e32 v7, v56, v46
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v7, -v7, v7, s0
	v_mul_f32_e32 v45, v54, v7
	v_mul_f32_e32 v47, v6, v7
	v_mul_f32_e32 v7, v53, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v45, v51, v46, -v45
	v_fma_f32 v47, v50, v46, -v47
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v7, v49, v46, -v7
	v_mul_f32_e32 v48, v45, v45
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v48, v47, v47
	v_fmac_f32_e32 v48, v7, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v46, 0x4f800000, v48
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v48
	v_cndmask_b32_e32 v46, v48, v46, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v48, v46
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v49, -1, v48
	v_add_nc_u32_e32 v50, 1, v48
	v_fma_f32 v51, -v49, v48, v46
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v55, -v50, v48, v46
	v_cmp_ge_f32_e64 s1, 0, v51
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v48, v48, v49, s1
	v_cmp_lt_f32_e64 s1, 0, v55
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v48, v48, v50, s1
	v_mul_f32_e32 v49, 0x37800000, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v48, v48, v49, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v46, 0x260
	v_cndmask_b32_e32 v46, v48, v46, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f32 v48, null, v46, v46, 1.0
	v_rcp_f32_e32 v49, v48
	s_waitcnt_depctr 0xfff
	v_fma_f32 v50, -v48, v49, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v49, v50, v49
	v_div_scale_f32 v50, vcc_lo, 1.0, v46, 1.0
	v_mul_f32_e32 v51, v50, v49
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v55, -v48, v51, v50
	v_fmac_f32_e32 v51, v55, v49
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v48, -v48, v51, v50
	v_div_fmas_f32 v48, v48, v49, v51
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v46, v48, v46, 1.0
                                        ; implicit-def: $vgpr48
	v_mul_f32_e32 v45, v45, v46
	v_mul_f32_e32 v50, v47, v46
	v_mul_f32_e32 v49, v7, v46
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v47, v54, v45
	v_fmac_f32_e32 v47, v6, v50
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v47, v53, v49
	v_cndmask_b32_e64 v7, v47, -v56, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v7, 1.0, v7
	v_mul_f32_e32 v46, 0x3f75c28f, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v46, v7, v46
	v_mul_f32_e32 v46, v7, v46
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v46, v7, v46
	v_fmaak_f32 v51, v7, v46, 0x3d23d70a
	v_lshrrev_b32_e32 v7, 6, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f64_f32_e32 v[46:47], v51
	v_xor_b32_e32 v7, v7, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[46:47], v[46:47], 0.5, 0x3fd00000
	v_cvt_f32_f64_e32 v53, v[46:47]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v46, 17, v7
                                        ; implicit-def: $vgpr47
	v_xor_b32_e32 v7, v46, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v46, 9, v7
	v_xor_b32_e32 v7, v46, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_or_b32 v46, 0x7fffff, v7, 1.0
	v_add_f32_e32 v46, -1.0, v46
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_nlt_f32_e32 vcc_lo, v46, v53
                                        ; implicit-def: $vgpr46
	s_and_saveexec_b32 s0, vcc_lo
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_33
; %bb.32:                               ;   in Loop: Header=BB0_13 Depth=4
	v_dual_sub_f32 v23, 1.0, v51 :: v_dual_sub_f32 v24, 1.0, v53
                                        ; implicit-def: $vgpr53
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v46, null, v24, v24, v23
	v_div_scale_f32 v51, vcc_lo, v23, v24, v23
	v_rcp_f32_e32 v47, v46
	s_waitcnt_depctr 0xfff
	v_fma_f32 v48, -v46, v47, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v47, v48, v47
	v_mul_f32_e32 v48, v51, v47
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v52, -v46, v48, v51
	v_fmac_f32_e32 v48, v52, v47
                                        ; implicit-def: $vgpr52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v46, -v46, v48, v51
                                        ; implicit-def: $vgpr51
	v_div_fmas_f32 v46, v46, v47, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v23, v46, v24, v23
                                        ; implicit-def: $vgpr24
	v_mul_f32_e32 v47, v23, v0
	v_mul_f32_e32 v48, v23, v1
	v_mul_f32_e32 v46, v23, v2
                                        ; implicit-def: $vgpr23
.LBB0_33:                               ; %Flow
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB0_35
; %bb.34:                               ;   in Loop: Header=BB0_13 Depth=4
	v_div_scale_f32 v45, null, v53, v53, v51
	v_div_scale_f32 v48, vcc_lo, v51, v53, v51
	v_mov_b32_e32 v50, v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v46, v45
	s_waitcnt_depctr 0xfff
	v_fma_f32 v47, -v45, v46, 1.0
	v_fmac_f32_e32 v46, v47, v46
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v47, v48, v46
	v_fma_f32 v49, -v45, v47, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v47, v49, v46
	v_mov_b32_e32 v49, v24
	v_fma_f32 v45, -v45, v47, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v45, v45, v46, v47
	v_div_fixup_f32 v46, v45, v53, v51
	v_mov_b32_e32 v45, v52
	s_delay_alu instid0(VALU_DEP_2)
	v_mul_f32_e32 v47, v46, v0
	v_mul_f32_e32 v48, v46, v1
	v_mul_f32_e32 v46, v46, v2
.LBB0_35:                               ; %Flow255
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s0
	v_dual_mov_b32 v24, v49 :: v_dual_mov_b32 v23, v50
	v_dual_mov_b32 v52, v45 :: v_dual_mov_b32 v45, v7
.LBB0_36:                               ; %Flow256
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s28
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s2, s2, exec_lo
.LBB0_37:                               ; %Flow262
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s27
	s_and_saveexec_b32 s27, s2
	s_cbranch_execz .LBB0_47
; %bb.38:                               ;   in Loop: Header=BB0_13 Depth=4
	v_lshrrev_b32_e32 v7, 6, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v7, v17
	v_lshlrev_b32_e32 v17, 17, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v17, v7
	v_lshrrev_b32_e32 v17, 9, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v24, v17, v7
	v_and_or_b32 v7, 0x7fffff, v24, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v7, -1.0, v7
	v_cvt_f64_f32_e32 v[45:46], v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[45:46], v[45:46], s[12:13]
	v_cvt_f32_f64_e32 v23, v[45:46]
                                        ; implicit-def: $vgpr46
                                        ; implicit-def: $vgpr45
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_and_b32_e32 v7, 0x7fffffff, v23
	v_cmp_ngt_f32_e64 s2, 0x48000000, |v23|
	v_lshrrev_b32_e32 v17, 23, v7
	v_and_or_b32 v47, 0x7fffff, v7, s19
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v48, 0xffffff88, v17
	s_and_saveexec_b32 s0, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s28, exec_lo, s0
	s_cbranch_execz .LBB0_40
; %bb.39:                               ;   in Loop: Header=BB0_13 Depth=4
	v_mad_u64_u32 v[45:46], null, 0xfe5163ab, v47, 0
	v_cmp_lt_u32_e32 vcc_lo, 63, v48
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v17, v46
	v_cndmask_b32_e64 v46, 0, 0xffffffc0, vcc_lo
	v_add_nc_u32_e32 v46, v46, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_u32_e64 s0, 31, v46
	v_cndmask_b32_e64 v55, 0, 0xffffffe0, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v46, v55, v46
	v_mad_u64_u32 v[49:50], null, 0x3c439041, v47, v[17:18]
	v_cmp_lt_u32_e64 s1, 31, v46
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v17, v50
	v_mad_u64_u32 v[50:51], null, 0xdb629599, v47, v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v17, v51
	v_cndmask_b32_e32 v45, v50, v45, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[51:52], null, 0xf534ddc0, v47, v[17:18]
	v_mov_b32_e32 v17, v52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[52:53], null, 0xfc2757d1, v47, v[17:18]
	v_mov_b32_e32 v17, v53
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[53:54], null, 0x4e441529, v47, v[17:18]
	v_dual_mov_b32 v17, v54 :: v_dual_cndmask_b32 v56, v53, v51
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[54:55], null, 0xa2f9836e, v47, v[17:18]
	v_cndmask_b32_e64 v17, 0, 0xffffffe0, s1
	v_dual_cndmask_b32 v54, v54, v52 :: v_dual_add_nc_u32 v17, v17, v46
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_cndmask_b32 v53, v55, v53 :: v_dual_cndmask_b32 v52, v52, v50
	v_cndmask_b32_e32 v46, v51, v49, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, 0, v17
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v49, v54, v56, s0
	v_cndmask_b32_e64 v51, v53, v54, s0
	v_cndmask_b32_e64 v53, v56, v52, s0
	v_sub_nc_u32_e32 v54, 32, v17
	v_cndmask_b32_e64 v52, v52, v46, s0
	v_cndmask_b32_e64 v45, v46, v45, s0
	v_cndmask_b32_e64 v51, v51, v49, s1
	v_cndmask_b32_e64 v49, v49, v53, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v53, v53, v52, s1
	v_cndmask_b32_e64 v45, v52, v45, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v55, v51, v49, v54
	v_alignbit_b32 v50, v49, v53, v54
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v52, v53, v45, v54
	v_dual_cndmask_b32 v17, v55, v51 :: v_dual_cndmask_b32 v46, v50, v49
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v52, v52, v53, vcc_lo
	v_bfe_u32 v49, v17, 29, 1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v50, v17, v46, 30
	v_alignbit_b32 v46, v46, v52, 30
	v_alignbit_b32 v45, v52, v45, 30
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v51, 0, v49
	v_xor_b32_e32 v50, v50, v51
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v46, v46, v51
	v_xor_b32_e32 v45, v45, v51
	v_lshrrev_b32_e32 v51, 29, v17
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_clz_i32_u32_e32 v53, v50
	v_lshlrev_b32_e32 v51, 31, v51
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_min_u32_e32 v53, 32, v53
	v_or_b32_e32 v54, 0.5, v51
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v52, 31, v53
	v_lshlrev_b32_e32 v55, 23, v53
	v_alignbit_b32 v50, v50, v46, v52
	v_alignbit_b32 v45, v46, v45, v52
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v54, v54, v55
	v_alignbit_b32 v46, v50, v45, 9
	v_lshrrev_b32_e32 v50, 9, v50
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v52, v46
	v_min_u32_e32 v52, 32, v52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v56, 31, v52
	v_alignbit_b32 v45, v46, v45, v56
	v_or_b32_e32 v46, v50, v54
	v_add_lshl_u32 v50, v52, v53, 23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshrrev_b32_e32 v45, 9, v45
	v_mul_f32_e32 v52, 0x3fc90fda, v46
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v45, v45, v50
	v_fma_f32 v50, 0x3fc90fda, v46, -v52
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v45, 0x33000000, v45
	v_fmac_f32_e32 v50, 0x33a22168, v46
	v_lshrrev_b32_e32 v17, 30, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v45, v45, v51
	v_add_nc_u32_e32 v46, v49, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v50, 0x3fc90fda, v45
	v_add_f32_e32 v45, v52, v50
.LBB0_40:                               ; %Flow258
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_saveexec_b32 s0, s28
	v_mul_f32_e64 v17, 0x3f22f983, |v23|
	s_delay_alu instid0(VALU_DEP_1)
	v_rndne_f32_e32 v50, v17
	s_xor_b32 exec_lo, exec_lo, s0
; %bb.41:                               ;   in Loop: Header=BB0_13 Depth=4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v45, 0xbfc90fda, v50, |v23|
	v_cvt_i32_f32_e32 v46, v50
	v_fmac_f32_e32 v45, 0xb3a22168, v50
	s_delay_alu instid0(VALU_DEP_1)
	v_fmac_f32_e32 v45, 0xa7c234c4, v50
; %bb.42:                               ; %_ZL3cosf.exit.i
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s0
                                        ; implicit-def: $vgpr49
                                        ; implicit-def: $vgpr17
	s_and_saveexec_b32 s0, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s2, exec_lo, s0
	s_cbranch_execz .LBB0_44
; %bb.43:                               ;   in Loop: Header=BB0_13 Depth=4
	v_mad_u64_u32 v[49:50], null, 0xfe5163ab, v47, 0
	v_cmp_lt_u32_e32 vcc_lo, 63, v48
	v_cndmask_b32_e64 v55, 0, 0xffffffc0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v17, v50 :: v_dual_add_nc_u32 v48, v55, v48
	v_mad_u64_u32 v[50:51], null, 0x3c439041, v47, v[17:18]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_u32_e64 s0, 31, v48
	v_cndmask_b32_e64 v56, 0, 0xffffffe0, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v17, v51 :: v_dual_add_nc_u32 v48, v56, v48
	v_mad_u64_u32 v[51:52], null, 0xdb629599, v47, v[17:18]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_lt_u32_e64 s1, 31, v48
	v_mov_b32_e32 v17, v52
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v49, v51, v49, vcc_lo
	v_mad_u64_u32 v[52:53], null, 0xf534ddc0, v47, v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v17, v53
	v_mad_u64_u32 v[53:54], null, 0xfc2757d1, v47, v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v17, v54
	v_mad_u64_u32 v[54:55], null, 0x4e441529, v47, v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mov_b32_e32 v17, v55
	v_mad_u64_u32 v[55:56], null, 0xa2f9836e, v47, v[17:18]
	v_cndmask_b32_e64 v17, 0, 0xffffffe0, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v47, v54, v52, vcc_lo
	v_dual_cndmask_b32 v48, v52, v50 :: v_dual_add_nc_u32 v17, v17, v48
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_cndmask_b32 v55, v55, v53 :: v_dual_cndmask_b32 v54, v56, v54
	v_cndmask_b32_e32 v53, v53, v51, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, 0, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v50, v55, v47, s0
	v_cndmask_b32_e64 v52, v54, v55, s0
	s_delay_alu instid0(VALU_DEP_4)
	v_cndmask_b32_e64 v47, v47, v53, s0
	v_sub_nc_u32_e32 v54, 32, v17
	v_cndmask_b32_e64 v53, v53, v48, s0
	v_cndmask_b32_e64 v48, v48, v49, s0
	v_cndmask_b32_e64 v52, v52, v50, s1
	v_cndmask_b32_e64 v50, v50, v47, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v47, v47, v53, s1
	v_cndmask_b32_e64 v48, v53, v48, s1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v55, v52, v50, v54
	v_alignbit_b32 v51, v50, v47, v54
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v53, v47, v48, v54
	v_cndmask_b32_e32 v17, v55, v52, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v49, v51, v50, vcc_lo
	v_cndmask_b32_e32 v47, v53, v47, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_bfe_u32 v50, v17, 29, 1
	v_alignbit_b32 v51, v17, v49, 30
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_alignbit_b32 v49, v49, v47, 30
	v_alignbit_b32 v47, v47, v48, 30
	v_sub_nc_u32_e32 v52, 0, v50
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v51, v51, v52
	v_xor_b32_e32 v48, v49, v52
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v47, v47, v52
	v_clz_i32_u32_e32 v53, v51
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v53, 32, v53
	v_sub_nc_u32_e32 v49, 31, v53
	v_lshlrev_b32_e32 v55, 23, v53
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_alignbit_b32 v51, v51, v48, v49
	v_alignbit_b32 v47, v48, v47, v49
	v_lshrrev_b32_e32 v49, 29, v17
	v_alignbit_b32 v48, v51, v47, 9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b32_e32 v49, 31, v49
	v_lshrrev_b32_e32 v51, 9, v51
	v_clz_i32_u32_e32 v52, v48
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v54, 0.5, v49
	v_min_u32_e32 v52, 32, v52
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v54, v54, v55
	v_sub_nc_u32_e32 v56, 31, v52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_alignbit_b32 v47, v48, v47, v56
	v_or_b32_e32 v48, v51, v54
	v_add_lshl_u32 v51, v52, v53, 23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshrrev_b32_e32 v47, 9, v47
	v_mul_f32_e32 v52, 0x3fc90fda, v48
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v47, v47, v51
	v_fma_f32 v51, 0x3fc90fda, v48, -v52
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v47, 0x33000000, v47
	v_fmac_f32_e32 v51, 0x33a22168, v48
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v47, v47, v49
	v_fmac_f32_e32 v51, 0x3fc90fda, v47
	v_lshrrev_b32_e32 v47, 30, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v17, v52, v51
	v_add_nc_u32_e32 v49, v50, v47
                                        ; implicit-def: $vgpr50
.LBB0_44:                               ; %Flow257
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_and_not1_saveexec_b32 s0, s2
; %bb.45:                               ;   in Loop: Header=BB0_13 Depth=4
	v_fma_f32 v17, 0xbfc90fda, v50, |v23|
	v_cvt_i32_f32_e32 v49, v50
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v17, 0xb3a22168, v50
	v_fmac_f32_e32 v17, 0xa7c234c4, v50
; %bb.46:                               ; %_ZL3sinf.exit.i
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s0
	v_cvt_f64_f32_e64 v[47:48], |v6|
	v_lshrrev_b32_e32 v53, 6, v24
	v_mul_f32_e32 v50, 0, v5
	v_dual_mul_f32 v56, v45, v45 :: v_dual_and_b32 v57, 1, v46
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v24, v53, v24
	v_and_b32_e32 v53, 1, v49
	v_fmaak_f32 v59, s21, v56, 0xbab64f3b
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_cmp_eq_u32_e64 s2, 0, v57
	v_cmp_lt_f64_e32 vcc_lo, s[14:15], v[47:48]
	v_cndmask_b32_e64 v6, 1.0, 0, vcc_lo
	v_cndmask_b32_e64 v48, 0, 1.0, vcc_lo
	v_dual_mul_f32 v47, v6, v3 :: v_dual_mul_f32 v52, v48, v4
	v_fma_f32 v48, v48, v3, -v50
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v47, v4, 0, -v47
	v_fma_f32 v6, v6, v5, -v52
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v51, v47, v47
	v_fmac_f32_e32 v51, v48, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v51, v6, v6
	v_mul_f32_e32 v50, 0x4f800000, v51
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v51
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v50, v51, v50, vcc_lo
	v_sqrt_f32_e32 v51, v50
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v52, -1, v51
	v_add_nc_u32_e32 v54, 1, v51
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v55, -v52, v51, v50
	v_fma_f32 v58, -v54, v51, v50
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cmp_ge_f32_e64 s0, 0, v55
	v_xor_b32_e32 v55, v7, v23
	v_dual_fmaak_f32 v7, s20, v56, 0x3c0881c4 :: v_dual_lshlrev_b32 v46, 30, v46
	v_cndmask_b32_e64 v51, v51, v52, s0
	v_cmp_lt_f32_e64 s0, 0, v58
	v_mul_f32_e32 v52, v17, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v51, v51, v54, s0
	v_dual_mul_f32 v58, 0x37800000, v51 :: v_dual_lshlrev_b32 v49, 30, v49
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v51, v51, v58 :: v_dual_lshlrev_b32 v54, 17, v24
	v_xor_b32_e32 v24, v54, v24
	v_cmp_class_f32_e64 vcc_lo, v50, 0x260
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v49, 0x80000000, v49
	v_lshrrev_b32_e32 v60, 9, v24
	v_dual_cndmask_b32 v50, v51, v50 :: v_dual_fmaak_f32 v51, v56, v7, 0xbe2aaa9d
	v_and_b32_e32 v46, 0x80000000, v46
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v7, v60, v24
	v_fmaak_f32 v24, v56, v59, 0x3d2aabf7
	v_div_scale_f32 v59, null, v50, v50, 1.0
	v_dual_mul_f32 v51, v56, v51 :: v_dual_fmaak_f32 v54, s20, v52, 0x3c0881c4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_or_b32 v60, 0x7fffff, v7, 1.0
	v_rcp_f32_e32 v61, v59
	v_div_scale_f32 v62, vcc_lo, 1.0, v50, 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_fmac_f32 v45, v45, v51 :: v_dual_fmaak_f32 v54, v52, v54, 0xbe2aaa9d
	v_add_f32_e32 v60, -1.0, v60
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_mul_f32 v54, v52, v54 :: v_dual_mul_f32 v51, 0x4f800000, v60
	v_cmp_gt_f32_e64 s0, 0xf800000, v60
	v_dual_fmaak_f32 v58, s21, v52, 0xbab64f3b :: v_dual_fmac_f32 v17, v17, v54
	s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v54, -v59, v61, 1.0
	v_cndmask_b32_e64 v51, v60, v51, s0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v61, v54, v61
	v_sqrt_f32_e32 v54, v51
	v_fmaak_f32 v24, v56, v24, 0xbf000004
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_fma_f32 v24, v56, v24, 1.0
	s_waitcnt_depctr 0xfff
	v_dual_sub_f32 v56, 1.0, v60 :: v_dual_add_nc_u32 v57, 1, v54
	v_cndmask_b32_e64 v24, -v45, v24, s2
	v_cmp_eq_u32_e64 s2, 0, v53
	v_dual_mul_f32 v60, 0x4f800000, v56 :: v_dual_add_nc_u32 v53, -1, v54
	v_cmp_gt_f32_e64 s1, 0xf800000, v56
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v24, v46, v24
	v_fma_f32 v46, -v53, v54, v51
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v56, v56, v60, s1
	v_sqrt_f32_e32 v45, v56
	v_fmaak_f32 v58, v52, v58, 0x3d2aabf7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmaak_f32 v58, v52, v58, 0xbf000004
	v_fma_f32 v52, v52, v58, 1.0
	v_mul_f32_e32 v58, v62, v61
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v17, v52, v17, s2
	v_fma_f32 v52, -v59, v58, v62
	v_cmp_ge_f32_e64 s2, 0, v46
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor3_b32 v17, v55, v49, v17
	v_fmac_f32_e32 v58, v52, v61
	v_fma_f32 v49, -v57, v54, v51
	v_add_nc_u32_e32 v52, -1, v45
	v_add_nc_u32_e32 v55, 1, v45
	v_cndmask_b32_e64 v46, v54, v53, s2
	v_fma_f32 v59, -v59, v58, v62
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v54, -v52, v45, v56
	v_div_fmas_f32 v53, v59, v61, v58
	v_cmp_lt_f32_e32 vcc_lo, 0, v49
	v_fma_f32 v58, -v55, v45, v56
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_div_fixup_f32 v49, v53, v50, 1.0
	v_cndmask_b32_e32 v46, v46, v57, vcc_lo
	v_cmp_ge_f32_e32 vcc_lo, 0, v54
	v_cmp_lt_f32_e64 s2, 0, v58
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v47, v47, v49
	v_dual_mul_f32 v48, v48, v49 :: v_dual_cndmask_b32 v45, v45, v52
	v_cmp_class_f32_e64 vcc_lo, v23, 0x1f8
	v_mul_f32_e32 v50, v3, v47
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v45, v45, v55, s2
	v_cndmask_b32_e32 v23, 0x7fc00000, v24, vcc_lo
	v_cndmask_b32_e32 v17, 0x7fc00000, v17, vcc_lo
	v_mul_f32_e32 v24, 0x37800000, v46
	v_cmp_class_f32_e64 vcc_lo, v51, 0x260
	v_cndmask_b32_e64 v24, v46, v24, s0
	v_mul_f32_e32 v46, 0x37800000, v45
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v24, v24, v51, vcc_lo
	v_cndmask_b32_e64 v45, v45, v46, s1
	v_mul_f32_e32 v46, v47, v23
	v_dual_mul_f32 v6, v6, v49 :: v_dual_mul_f32 v51, v5, v48
	v_cmp_class_f32_e64 vcc_lo, v56, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v49, v4, v6
	v_fma_f32 v47, v4, v47, -v51
	v_cndmask_b32_e32 v45, v45, v56, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v49, v3, v48, -v49
	v_mul_f32_e32 v49, v49, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v49, v24, v49
	v_fmac_f32_e32 v49, v24, v46
	v_fma_f32 v50, v5, v6, -v50
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_mov_b32 v46, v2 :: v_dual_fmac_f32 v49, v5, v45
	v_mul_f32_e32 v50, v50, v17
	v_dual_mul_f32 v17, v47, v17 :: v_dual_mul_f32 v6, v6, v23
	v_mov_b32_e32 v47, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v50, v24, v50
	v_mul_f32_e32 v17, v24, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v17, v24, v6
	v_dual_mul_f32 v48, v48, v23 :: v_dual_fmac_f32 v17, v3, v45
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v50, v24, v48
	v_mov_b32_e32 v48, v1
	v_dual_fmac_f32 v50, v4, v45 :: v_dual_mov_b32 v45, v7
	v_mul_f32_e32 v4, v49, v49
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v50, v50
	v_fmac_f32_e32 v4, v17, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v3, 0x4f800000, v4
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v4
	v_cndmask_b32_e32 v3, v4, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v5, -1, v4
	v_add_nc_u32_e32 v6, 1, v4
	v_fma_f32 v23, -v5, v4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v24, -v6, v4, v3
	v_cmp_ge_f32_e64 s0, 0, v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v4, v4, v5, s0
	v_cmp_lt_f32_e64 s0, 0, v24
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v4, v4, v6, s0
	v_mul_f32_e32 v5, 0x37800000, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v3, 0x260
	v_cndmask_b32_e32 v3, v4, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v4, null, v3, v3, 1.0
	v_div_scale_f32 v23, vcc_lo, 1.0, v3, 1.0
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v4, v5, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v6, v5
	v_mul_f32_e32 v6, v23, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v24, -v4, v6, v23
	v_fmac_f32_e32 v6, v24, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v4, -v4, v6, v23
	v_div_fmas_f32 v4, v4, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v3, v4, v3, 1.0
	v_mul_f32_e32 v23, v50, v3
	v_mul_f32_e32 v52, v49, v3
	v_mul_f32_e32 v24, v17, v3
.LBB0_47:                               ; %Flow263
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s27
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_mov_b32 v49, v24 :: v_dual_mov_b32 v50, v23
	v_mov_b32_e32 v51, v52
	v_mov_b32_e32 v17, v7
	s_xor_b32 s2, exec_lo, -1
.LBB0_48:                               ; %Flow267
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s26
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v0, v17
	s_or_not1_b32 s0, s2, exec_lo
.LBB0_49:                               ; %Flow264
                                        ;   in Loop: Header=BB0_13 Depth=4
	s_or_b32 exec_lo, exec_lo, s3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, exec_lo, s0
	s_or_b32 s25, s0, s25
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s25
	s_cbranch_execnz .LBB0_13
; %bb.50:                               ; %_Z8radianceIPK6SphereE3VecRK3RayiT_R3RNG.exit
                                        ;   in Loop: Header=BB0_4 Depth=3
	s_or_b32 exec_lo, exec_lo, s25
	v_fmac_f32_e32 v36, v40, v26
	v_fmac_f32_e32 v35, v38, v26
	v_fmac_f32_e32 v34, v39, v26
	s_add_i32 s24, s24, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s24, s10
	s_cbranch_scc0 .LBB0_4
	s_branch .LBB0_52
.LBB0_51:                               ;   in Loop: Header=BB0_2 Depth=2
	v_dual_mov_b32 v35, 0 :: v_dual_mov_b32 v36, 0
.LBB0_52:                               ; %Flow271
                                        ;   in Loop: Header=BB0_2 Depth=2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_nlt_f32_e32 vcc_lo, 1.0, v36
	s_xor_b32 s0, s23, -1
	s_mov_b32 s16, 0
	s_mov_b32 s17, 0x3ff80000
	s_mov_b32 s23, 0
	v_cndmask_b32_e32 v1, 1.0, v36, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 1.0, v35
	v_cndmask_b32_e32 v2, 1.0, v35, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 1.0, v34
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_dual_mul_f32 v1, 0x3e800000, v1 :: v_dual_mul_f32 v2, 0x3e800000, v2
	v_cndmask_b32_e32 v3, 1.0, v34, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0, v36
	v_cndmask_b32_e32 v1, 0, v1, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0, v35
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_dual_add_f32 v8, v1, v8 :: v_dual_mul_f32 v3, 0x3e800000, v3
	v_cndmask_b32_e32 v2, 0, v2, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0, v34
	v_add_f32_e32 v9, v2, v9
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v3, 0, v3, vcc_lo
	s_and_b32 vcc_lo, exec_lo, s0
	v_add_f32_e32 v10, v3, v10
	global_store_b96 v[15:16], v[8:10], off
	s_cbranch_vccz .LBB0_2
; %bb.53:                               ;   in Loop: Header=BB0_1 Depth=1
	s_mov_b32 s8, 0
	s_mov_b32 s9, 0x3ff80000
	s_and_b32 vcc_lo, exec_lo, s22
	s_mov_b32 s0, 0
	s_cbranch_vccz .LBB0_1
; %bb.54:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8raytraceP3VecPK6Sphereiii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 63
		.amdhsa_next_free_sgpr 29
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z8raytraceP3VecPK6Sphereiii, .Lfunc_end0-_Z8raytraceP3VecPK6Sphereiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 7716
; NumSgprs: 31
; NumVgprs: 63
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 7
; NumSGPRsForWavesPerEU: 31
; NumVGPRsForWavesPerEU: 63
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8raytraceP3VecPK6Sphereiii
    .private_segment_fixed_size: 0
    .sgpr_count:     31
    .sgpr_spill_count: 0
    .symbol:         _Z8raytraceP3VecPK6Sphereiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     63
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
