#include <iostream>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <vector>

#define BLOCK_SIZE 16
#define NX 512
#define NY 512
#define NZ 512
#define NT 32
#define NDIR 19

__constant__ float c[NDIR][3];
__constant__ float t[NDIR];

// Lattice Boltzmann kernel
__global__ void lbm_kernel(float *f, float *f_new, float *rho, float *ux, float *uy, float *uz, float omega) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;

    if (x >= NX || y >= NY || z >= NZ) return;

    // Calculate fluid density
    float rho_sum = 0.0f;
    for (int i = 0; i < NDIR; ++i) {
        rho_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x];
    }
    rho[z * NX * NY + y * NX + x] = rho_sum;

    // Calculate fluid velocity
    float ux_sum = 0.0f, uy_sum = 0.0f, uz_sum = 0.0f;
    for (int i = 0; i < NDIR; ++i) {
        ux_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x] * c[i][0];
        uy_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x] * c[i][1];
        uz_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x] * c[i][2];
    }
    ux[z * NX * NY + y * NX + x] = ux_sum / rho_sum;
    uy[z * NX * NY + y * NX + x] = uy_sum / rho_sum;
    uz[z * NX * NY + y * NX + x] = uz_sum / rho_sum;

    // Streaming step
    for (int i = 0; i < NDIR; ++i) {
        int xn = x + c[i][0];
        int yn = y + c[i][1];
        int zn = z + c[i][2];

        // Periodic boundary conditions
        xn = (xn + NX) % NX;
        yn = (yn + NY) % NY;
        zn = (zn + NZ) % NZ;

        f_new[i * NX * NY * NZ + zn * NX * NY + yn * NX + xn] = f[i * NX * NY * NZ + z * NX * NY + y * NX + x];
    }
}

// Collision step kernel
__global__ void collision_kernel(float *f, float *f_new, float *rho, float *ux, float *uy, float *uz, float omega) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;

    if (x >= NX || y >= NY || z >= NZ) return;

    float u_dot_c = 0.0f;
    float u_sq = ux[z * NX * NY + y * NX + x] * ux[z * NX * NY + y * NX + x] +
                 uy[z * NX * NY + y * NX + x] * uy[z * NX * NY + y * NX + x] +
                 uz[z * NX * NY + y * NX + x] * uz[z * NX * NY + y * NX + x];

    for (int i = 0; i < NDIR; ++i) {
        u_dot_c = ux[z * NX * NY + y * NX + x] * c[i][0] +
                  uy[z * NX * NY + y * NX + x] * c[i][1] +
                  uz[z * NX * NY + y * NX + x] * c[i][2];

        float feq = rho[z * NX * NY + y * NX + x] * t[i] *
                    (1.0f + 3.0f * u_dot_c + 4.5f * u_dot_c * u_dot_c - 1.5f * u_sq);

        f[i * NX * NY * NZ + z * NX * NY + y * NX + x] =
            f_new[i * NX * NY * NZ + z * NX * NY + y * NX + x] -
            omega * (f_new[i * NX * NY * NZ + z * NX * NY + y * NX + x] - feq);
    }
}

int main() {
    // Initialize direction vectors and weights
    float c_host[NDIR][3] = {
        {0, 0, 0}, {1, 0, 0}, {0, 1, 0}, {0, 0, 1}, {-1, 0, 0}, {0, -1, 0}, {0, 0, -1},
        {1, 1, 0}, {1, -1, 0}, {-1, 1, 0}, {-1, -1, 0}, {1, 0, 1}, {1, 0, -1}, {-1, 0, 1},
        {-1, 0, -1}, {0, 1, 1}, {0, 1, -1}, {0, -1, 1}, {0, -1, -1}
    };
    float t_host[NDIR] = {
        1.0f / 3.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f,
        1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f,
        1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f
    };

    cudaMemcpyToSymbol(c, c_host, sizeof(c_host));
    cudaMemcpyToSymbol(t, t_host, sizeof(t_host));

    // Allocate memory
    float *f, *f_new, *rho, *ux, *uy, *uz;
    cudaMalloc(&f, NDIR * NX * NY * NZ * sizeof(float));
    cudaMalloc(&f_new, NDIR * NX * NY * NZ * sizeof(float));
    cudaMalloc(&rho, NX * NY * NZ * sizeof(float));
    cudaMalloc(&ux, NX * NY * NZ * sizeof(float));
    cudaMalloc(&uy, NX * NY * NZ * sizeof(float));
    cudaMalloc(&uz, NX * NY * NZ * sizeof(float));

    // Initialize distribution functions
    cudaMemset(f, 0, NDIR * NX * NY * NZ * sizeof(float));
    cudaMemset(f_new, 0, NDIR * NX * NY * NZ * sizeof(float));

    // Set initial conditions (e.g., equilibrium distribution)
    // This is a simplified initialization for demonstration purposes
    float omega = 1.8f; // Relaxation parameter

    dim3 blockSize(BLOCK_SIZE, BLOCK_SIZE, BLOCK_SIZE);
    dim3 gridSize((NX + BLOCK_SIZE - 1) / BLOCK_SIZE, (NY + BLOCK_SIZE - 1) / BLOCK_SIZE, (NZ + BLOCK_SIZE - 1) / BLOCK_SIZE);

    for (int t = 0; t < NT; ++t) {
        lbm_kernel<<<gridSize, blockSize>>>(f, f_new, rho, ux, uy, uz, omega);
        cudaDeviceSynchronize();
        collision_kernel<<<gridSize, blockSize>>>(f, f_new, rho, ux, uy, uz, omega);
        cudaDeviceSynchronize();
    }

    // Free memory
    cudaFree(f);
    cudaFree(f_new);
    cudaFree(rho);
    cudaFree(ux);
    cudaFree(uy);
    cudaFree(uz);

    return 0;
}
	.file	"tmpxft_0011534f_00000000-6_cuda_code_077162.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4375:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4375:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
	.type	_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f, @function
_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f:
.LFB4397:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	72(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z10lbm_kernelPfS_S_S_S_S_f(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 224
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L2:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4397:
	.size	_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f, .-_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
	.globl	_Z10lbm_kernelPfS_S_S_S_S_f
	.type	_Z10lbm_kernelPfS_S_S_S_S_f, @function
_Z10lbm_kernelPfS_S_S_S_S_f:
.LFB4398:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
	.cfi_endproc
.LFE4398:
	.size	_Z10lbm_kernelPfS_S_S_S_S_f, .-_Z10lbm_kernelPfS_S_S_S_S_f
	.globl	_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
	.type	_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f, @function
_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f:
.LFB4399:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	%rdi, 56(%rsp)
	leaq	80(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	92(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	64(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	72(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 88(%rsp)
	movq	%rax, 128(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 80(%rsp)
	movq	%rax, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	72(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z16collision_kernelPfS_S_S_S_S_f(%rip), %rdi
	pushq	72(%rsp)
	.cfi_def_cfa_offset 224
	movq	108(%rsp), %rcx
	movl	116(%rsp), %r8d
	movq	96(%rsp), %rsi
	movl	104(%rsp), %edx
	leaq	144(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L8:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4399:
	.size	_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f, .-_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
	.globl	_Z16collision_kernelPfS_S_S_S_S_f
	.type	_Z16collision_kernelPfS_S_S_S_S_f, @function
_Z16collision_kernelPfS_S_S_S_S_f:
.LFB4400:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
	.cfi_endproc
.LFE4400:
	.size	_Z16collision_kernelPfS_S_S_S_S_f, .-_Z16collision_kernelPfS_S_S_S_S_f
	.section	.rodata
	.align 32
.LC0:
	.long	0
	.long	0
	.long	0
	.long	1065353216
	.long	0
	.long	0
	.long	0
	.long	1065353216
	.long	0
	.long	0
	.long	0
	.long	1065353216
	.long	-1082130432
	.long	0
	.long	0
	.long	0
	.long	-1082130432
	.long	0
	.long	0
	.long	0
	.long	-1082130432
	.long	1065353216
	.long	1065353216
	.long	0
	.long	1065353216
	.long	-1082130432
	.long	0
	.long	-1082130432
	.long	1065353216
	.long	0
	.long	-1082130432
	.long	-1082130432
	.long	0
	.long	1065353216
	.long	0
	.long	1065353216
	.long	1065353216
	.long	0
	.long	-1082130432
	.long	-1082130432
	.long	0
	.long	1065353216
	.long	-1082130432
	.long	0
	.long	-1082130432
	.long	0
	.long	1065353216
	.long	1065353216
	.long	0
	.long	1065353216
	.long	-1082130432
	.long	0
	.long	-1082130432
	.long	1065353216
	.long	0
	.long	-1082130432
	.long	-1082130432
	.align 32
.LC1:
	.long	1051372203
	.long	1029934649
	.long	1029934649
	.long	1029934649
	.long	1029934649
	.long	1029934649
	.long	1029934649
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.long	1021546041
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4372:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$57, %ecx
	movl	$228, %edx
	leaq	.LC0(%rip), %rsi
	movabsq	$-6979321856, %rbx
	movl	$1, %r8d
	subq	$384, %rsp
	.cfi_def_cfa_offset 400
	movq	%fs:40, %rax
	movq	%rax, 376(%rsp)
	xorl	%eax, %eax
	leaq	148(%rsp), %rdi
	rep movsl
	leaq	72(%rsp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$19, %ecx
	rep movsl
	leaq	148(%rsp), %rsi
	leaq	_ZL1c(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$76, %edx
	leaq	72(%rsp), %rsi
	leaq	_ZL1t(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	call	cudaMalloc@PLT
	movq	%rbx, %rsi
	leaq	8(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	16(%rsp), %rdi
	movl	$536870912, %esi
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movl	$536870912, %esi
	call	cudaMalloc@PLT
	leaq	32(%rsp), %rdi
	movl	$536870912, %esi
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movl	$536870912, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movq	%rbx, %rdx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	xorl	%esi, %esi
	movl	$32, %ebx
	call	cudaMemset@PLT
	movl	$16, 56(%rsp)
	movabsq	$68719476752, %rax
	movq	%rax, 48(%rsp)
	movabsq	$137438953504, %rax
	movq	%rax, 60(%rsp)
	movl	$32, 68(%rsp)
.L16:
	movl	56(%rsp), %ecx
	movq	48(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	60(%rsp), %rdi
	movl	68(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L14
	movss	.LC2(%rip), %xmm0
	movq	40(%rsp), %r9
	movq	32(%rsp), %r8
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	call	_Z41__device_stub__Z10lbm_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
.L14:
	call	cudaDeviceSynchronize@PLT
	movl	56(%rsp), %ecx
	movl	68(%rsp), %esi
	xorl	%r9d, %r9d
	movq	48(%rsp), %rdx
	movq	60(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	movss	.LC2(%rip), %xmm0
	movq	40(%rsp), %r9
	movq	32(%rsp), %r8
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	call	_Z47__device_stub__Z16collision_kernelPfS_S_S_S_S_fPfS_S_S_S_S_f
.L15:
	call	cudaDeviceSynchronize@PLT
	decl	%ebx
	jne	.L16
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	376(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$384, %rsp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4372:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC5:
	.string	"_Z16collision_kernelPfS_S_S_S_S_f"
.LC6:
	.string	"_Z10lbm_kernelPfS_S_S_S_S_f"
.LC7:
	.string	"c"
.LC8:
	.string	"t"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4402:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z16collision_kernelPfS_S_S_S_S_f(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10lbm_kernelPfS_S_S_S_S_f(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	movl	$228, %r9d
	leaq	_ZL1c(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC8(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$76, %r9d
	leaq	_ZL1t(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4402:
	.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
	.local	_ZL1t
	.comm	_ZL1t,76,32
	.local	_ZL1c
	.comm	_ZL1c,228,32
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC2:
	.long	1072064102
	.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:

#include <iostream>
#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
#include <vector>

#define BLOCK_SIZE 16
#define NX 512
#define NY 512
#define NZ 512
#define NT 32
#define NDIR 19

__constant__ float c[NDIR][3];
__constant__ float t[NDIR];

// Lattice Boltzmann kernel
__global__ void lbm_kernel(float *f, float *f_new, float *rho, float *ux, float *uy, float *uz, float omega) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;

    if (x >= NX || y >= NY || z >= NZ) return;

    // Calculate fluid density
    float rho_sum = 0.0f;
    for (int i = 0; i < NDIR; ++i) {
        rho_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x];
    }
    rho[z * NX * NY + y * NX + x] = rho_sum;

    // Calculate fluid velocity
    float ux_sum = 0.0f, uy_sum = 0.0f, uz_sum = 0.0f;
    for (int i = 0; i < NDIR; ++i) {
        ux_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x] * c[i][0];
        uy_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x] * c[i][1];
        uz_sum += f[i * NX * NY * NZ + z * NX * NY + y * NX + x] * c[i][2];
    }
    ux[z * NX * NY + y * NX + x] = ux_sum / rho_sum;
    uy[z * NX * NY + y * NX + x] = uy_sum / rho_sum;
    uz[z * NX * NY + y * NX + x] = uz_sum / rho_sum;

    // Streaming step
    for (int i = 0; i < NDIR; ++i) {
        int xn = x + c[i][0];
        int yn = y + c[i][1];
        int zn = z + c[i][2];

        // Periodic boundary conditions
        xn = (xn + NX) % NX;
        yn = (yn + NY) % NY;
        zn = (zn + NZ) % NZ;

        f_new[i * NX * NY * NZ + zn * NX * NY + yn * NX + xn] = f[i * NX * NY * NZ + z * NX * NY + y * NX + x];
    }
}

// Collision step kernel
__global__ void collision_kernel(float *f, float *f_new, float *rho, float *ux, float *uy, float *uz, float omega) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;

    if (x >= NX || y >= NY || z >= NZ) return;

    float u_dot_c = 0.0f;
    float u_sq = ux[z * NX * NY + y * NX + x] * ux[z * NX * NY + y * NX + x] +
                 uy[z * NX * NY + y * NX + x] * uy[z * NX * NY + y * NX + x] +
                 uz[z * NX * NY + y * NX + x] * uz[z * NX * NY + y * NX + x];

    for (int i = 0; i < NDIR; ++i) {
        u_dot_c = ux[z * NX * NY + y * NX + x] * c[i][0] +
                  uy[z * NX * NY + y * NX + x] * c[i][1] +
                  uz[z * NX * NY + y * NX + x] * c[i][2];

        float feq = rho[z * NX * NY + y * NX + x] * t[i] *
                    (1.0f + 3.0f * u_dot_c + 4.5f * u_dot_c * u_dot_c - 1.5f * u_sq);

        f[i * NX * NY * NZ + z * NX * NY + y * NX + x] =
            f_new[i * NX * NY * NZ + z * NX * NY + y * NX + x] -
            omega * (f_new[i * NX * NY * NZ + z * NX * NY + y * NX + x] - feq);
    }
}

int main() {
    // Initialize direction vectors and weights
    float c_host[NDIR][3] = {
        {0, 0, 0}, {1, 0, 0}, {0, 1, 0}, {0, 0, 1}, {-1, 0, 0}, {0, -1, 0}, {0, 0, -1},
        {1, 1, 0}, {1, -1, 0}, {-1, 1, 0}, {-1, -1, 0}, {1, 0, 1}, {1, 0, -1}, {-1, 0, 1},
        {-1, 0, -1}, {0, 1, 1}, {0, 1, -1}, {0, -1, 1}, {0, -1, -1}
    };
    float t_host[NDIR] = {
        1.0f / 3.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f, 1.0f / 18.0f,
        1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f,
        1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f, 1.0f / 36.0f
    };

    hipMemcpyToSymbol(HIP_SYMBOL(c), c_host, sizeof(c_host));
    hipMemcpyToSymbol(HIP_SYMBOL(t), t_host, sizeof(t_host));

    // Allocate memory
    float *f, *f_new, *rho, *ux, *uy, *uz;
    hipMalloc(&f, NDIR * NX * NY * NZ * sizeof(float));
    hipMalloc(&f_new, NDIR * NX * NY * NZ * sizeof(float));
    hipMalloc(&rho, NX * NY * NZ * sizeof(float));
    hipMalloc(&ux, NX * NY * NZ * sizeof(float));
    hipMalloc(&uy, NX * NY * NZ * sizeof(float));
    hipMalloc(&uz, NX * NY * NZ * sizeof(float));

    // Initialize distribution functions
    hipMemset(f, 0, NDIR * NX * NY * NZ * sizeof(float));
    hipMemset(f_new, 0, NDIR * NX * NY * NZ * sizeof(float));

    // Set initial conditions (e.g., equilibrium distribution)
    // This is a simplified initialization for demonstration purposes
    float omega = 1.8f; // Relaxation parameter

    dim3 blockSize(BLOCK_SIZE, BLOCK_SIZE, BLOCK_SIZE);
    dim3 gridSize((NX + BLOCK_SIZE - 1) / BLOCK_SIZE, (NY + BLOCK_SIZE - 1) / BLOCK_SIZE, (NZ + BLOCK_SIZE - 1) / BLOCK_SIZE);

    for (int t = 0; t < NT; ++t) {
        lbm_kernel<<<gridSize, blockSize>>>(f, f_new, rho, ux, uy, uz, omega);
        hipDeviceSynchronize();
        collision_kernel<<<gridSize, blockSize>>>(f, f_new, rho, ux, uy, uz, omega);
        hipDeviceSynchronize();
    }

    // Free memory
    hipFree(f);
    hipFree(f_new);
    hipFree(rho);
    hipFree(ux);
    hipFree(uy);
    hipFree(uz);

    return 0;
}
	.text
	.file	"cuda_code_077162.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f # -- Begin function _Z25__device_stub__lbm_kernelPfS_S_S_S_S_f
	.type	_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f,@function
_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f: # @_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	72(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	64(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	56(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	48(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	32(%rsp), %r8
	movq	%r9, (%r8)
	leaq	12(%rsp), %r9
	movss	%xmm0, (%r9)
	leaq	112(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	movq	%r9, 48(%rbx)
	leaq	96(%rsp), %r14
	leaq	80(%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	$_Z10lbm_kernelPfS_S_S_S_S_f, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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	_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f, .Lfunc_end0-_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f
	.cfi_endproc
                                        # -- End function
	.globl	_Z31__device_stub__collision_kernelPfS_S_S_S_S_f # -- Begin function _Z31__device_stub__collision_kernelPfS_S_S_S_S_f
	.type	_Z31__device_stub__collision_kernelPfS_S_S_S_S_f,@function
_Z31__device_stub__collision_kernelPfS_S_S_S_S_f: # @_Z31__device_stub__collision_kernelPfS_S_S_S_S_f
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	72(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	64(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	56(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	48(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	32(%rsp), %r8
	movq	%r9, (%r8)
	leaq	12(%rsp), %r9
	movss	%xmm0, (%r9)
	leaq	112(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	movq	%r9, 48(%rbx)
	leaq	96(%rsp), %r14
	leaq	80(%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	$_Z16collision_kernelPfS_S_S_S_S_f, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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_end1:
	.size	_Z31__device_stub__collision_kernelPfS_S_S_S_S_f, .Lfunc_end1-_Z31__device_stub__collision_kernelPfS_S_S_S_S_f
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI2_0:
	.long	0x3fe66666                      # float 1.79999995
	.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	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$360, %rsp                      # imm = 0x168
	.cfi_def_cfa_offset 416
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movabsq	$137438953504, %rbx             # imm = 0x2000000020
	movabsq	$68719476752, %r14              # imm = 0x1000000010
	leaq	128(%rsp), %r15
	movl	$.L__const.main.c_host, %esi
	movl	$228, %edx
	movq	%r15, %rdi
	callq	memcpy@PLT
	leaq	48(%rsp), %r12
	movl	$9, %ecx
	movl	$.L__const.main.t_host, %esi
	movq	%r12, %rdi
	rep;movsq (%rsi), %es:(%rdi)
	movl	$1021546041, 72(%r12)           # imm = 0x3CE38E39
	movl	$c, %edi
	movl	$228, %edx
	movq	%r15, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$t, %edi
	movl	$76, %edx
	movq	%r12, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	leaq	40(%rsp), %r12
	movabsq	$-6979321856, %r15              # imm = 0xFFFFFFFE60000000
	movq	%r12, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	32(%rsp), %r13
	movq	%r13, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	hipMalloc
	leaq	8(%rsp), %rdi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	hipMalloc
	movq	%rsp, %rdi
	movl	$536870912, %esi                # imm = 0x20000000
	callq	hipMalloc
	movq	(%r12), %rdi
	xorl	%esi, %esi
	movq	%r15, %rdx
	callq	hipMemset
	movq	(%r13), %rdi
	xorl	%esi, %esi
	movq	%r15, %rdx
	callq	hipMemset
	movl	$32, %ebp
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	movl	$32, %esi
	movq	%r14, %rdx
	movl	$16, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_3
# %bb.2:                                #   in Loop: Header=BB2_1 Depth=1
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	8(%rsp), %r8
	movq	(%rsp), %r9
	movss	.LCPI2_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	callq	_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f
.LBB2_3:                                #   in Loop: Header=BB2_1 Depth=1
	callq	hipDeviceSynchronize
	movq	%rbx, %rdi
	movl	$32, %esi
	movq	%r14, %rdx
	movl	$16, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_5
# %bb.4:                                #   in Loop: Header=BB2_1 Depth=1
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	8(%rsp), %r8
	movq	(%rsp), %r9
	movss	.LCPI2_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	callq	_Z31__device_stub__collision_kernelPfS_S_S_S_S_f
.LBB2_5:                                #   in Loop: Header=BB2_1 Depth=1
	callq	hipDeviceSynchronize
	decl	%ebp
	jne	.LBB2_1
# %bb.6:
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$360, %rsp                      # imm = 0x168
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	main, .Lfunc_end2-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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10lbm_kernelPfS_S_S_S_S_f, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z16collision_kernelPfS_S_S_S_S_f, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$c, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$228, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$t, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movl	$76, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	c,@object                       # @c
	.local	c
	.comm	c,228,16
	.type	t,@object                       # @t
	.local	t
	.comm	t,76,16
	.type	_Z10lbm_kernelPfS_S_S_S_S_f,@object # @_Z10lbm_kernelPfS_S_S_S_S_f
	.section	.rodata,"a",@progbits
	.globl	_Z10lbm_kernelPfS_S_S_S_S_f
	.p2align	3, 0x0
_Z10lbm_kernelPfS_S_S_S_S_f:
	.quad	_Z25__device_stub__lbm_kernelPfS_S_S_S_S_f
	.size	_Z10lbm_kernelPfS_S_S_S_S_f, 8

	.type	_Z16collision_kernelPfS_S_S_S_S_f,@object # @_Z16collision_kernelPfS_S_S_S_S_f
	.globl	_Z16collision_kernelPfS_S_S_S_S_f
	.p2align	3, 0x0
_Z16collision_kernelPfS_S_S_S_S_f:
	.quad	_Z31__device_stub__collision_kernelPfS_S_S_S_S_f
	.size	_Z16collision_kernelPfS_S_S_S_S_f, 8

	.type	.L__const.main.c_host,@object   # @__const.main.c_host
	.p2align	4, 0x0
.L__const.main.c_host:
	.zero	12
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0x3f800000                      # float 1
	.long	0xbf800000                      # float -1
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0x3f800000                      # float 1
	.long	0x00000000                      # float 0
	.long	0xbf800000                      # float -1
	.long	0xbf800000                      # float -1
	.size	.L__const.main.c_host, 228

	.type	.L__const.main.t_host,@object   # @__const.main.t_host
	.p2align	4, 0x0
.L__const.main.t_host:
	.long	0x3eaaaaab                      # float 0.333333343
	.long	0x3d638e39                      # float 0.055555556
	.long	0x3d638e39                      # float 0.055555556
	.long	0x3d638e39                      # float 0.055555556
	.long	0x3d638e39                      # float 0.055555556
	.long	0x3d638e39                      # float 0.055555556
	.long	0x3d638e39                      # float 0.055555556
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.long	0x3ce38e39                      # float 0.027777778
	.size	.L__const.main.t_host, 76

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z10lbm_kernelPfS_S_S_S_S_f"
	.size	.L__unnamed_1, 28

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z16collision_kernelPfS_S_S_S_S_f"
	.size	.L__unnamed_2, 34

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"c"
	.size	.L__unnamed_3, 2

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"t"
	.size	.L__unnamed_4, 2

	.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 _Z25__device_stub__lbm_kernelPfS_S_S_S_S_f
	.addrsig_sym _Z31__device_stub__collision_kernelPfS_S_S_S_S_f
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym c
	.addrsig_sym t
	.addrsig_sym _Z10lbm_kernelPfS_S_S_S_S_f
	.addrsig_sym _Z16collision_kernelPfS_S_S_S_S_f
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z10lbm_kernelPfS_S_S_S_S_f ; -- Begin function _Z10lbm_kernelPfS_S_S_S_S_f
	.globl	_Z10lbm_kernelPfS_S_S_S_S_f
	.p2align	8
	.type	_Z10lbm_kernelPfS_S_S_S_S_f,@function
_Z10lbm_kernelPfS_S_S_S_S_f:            ; @_Z10lbm_kernelPfS_S_S_S_S_f
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x44
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v2, v0, 10, 10
	v_bfe_u32 v5, v0, 20, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[3:4], null, s13, s2, v[1:2]
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, s14, s4, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, s3, v[5:6]
	v_max3_i32 v2, v3, v0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e32 0x200, v2
	s_cbranch_execz .LBB0_7
; %bb.1:                                ; %.preheader
	s_clause 0x1
	s_load_b256 s[4:11], s[0:1], 0x0
	s_load_b128 s[12:15], s[0:1], 0x20
	v_lshlrev_b32_e32 v6, 18, v1
	v_dual_mov_b32 v8, 0 :: v_dual_lshlrev_b32 v7, 9, v0
	s_mov_b32 s0, 19
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v2, v3, v6, v7
	v_mov_b32_e32 v4, v2
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_ashrrev_i32_e32 v5, 31, v4
	s_add_i32 s0, s0, -1
	s_cmp_eq_u32 s0, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[9:10], 2, v[4:5]
	v_add_nc_u32_e32 v4, 0x8000000, v4
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v9, vcc_lo, s4, v9
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v10, vcc_lo
	global_load_b32 v5, v[9:10], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v8, v8, v5
	s_cbranch_scc0 .LBB0_2
; %bb.3:
	v_add3_u32 v4, v7, v3, v6
	v_dual_mov_b32 v9, 0 :: v_dual_mov_b32 v6, v2
	v_dual_mov_b32 v10, 0 :: v_dual_mov_b32 v11, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v5, 31, v4
	s_mov_b64 s[0:1], 8
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v12, vcc_lo, s8, v4
	v_add_co_ci_u32_e32 v13, vcc_lo, s9, v5, vcc_lo
	global_store_b32 v[12:13], v8, off
.LBB0_4:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_3)
	v_ashrrev_i32_e32 v7, 31, v6
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, c@rel32@lo-4
	s_addc_u32 s3, s3, c@rel32@hi+4
	s_add_u32 s2, s0, s2
	s_addc_u32 s3, s1, s3
	s_add_u32 s0, s0, 12
	v_lshlrev_b64 v[12:13], 2, v[6:7]
	s_addc_u32 s1, s1, 0
	s_cmpk_eq_i32 s0, 0xec
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v12, vcc_lo, s4, v12
	v_add_co_ci_u32_e32 v13, vcc_lo, s5, v13, vcc_lo
	global_load_b32 v7, v[12:13], off
	s_clause 0x1
	s_load_b64 s[8:9], s[2:3], 0x0
	s_load_b32 s2, s[2:3], 0x8
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_dual_fmac_f32 v11, s8, v7 :: v_dual_add_nc_u32 v6, 0x8000000, v6
	v_fmac_f32_e32 v10, s9, v7
	v_fmac_f32_e32 v9, s2, v7
	s_cbranch_scc0 .LBB0_4
; %bb.5:
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_scale_f32 v14, null, v8, v8, v11
	v_div_scale_f32 v15, null, v8, v8, v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_scale_f32 v16, null, v8, v8, v9
	v_rcp_f32_e32 v17, v14
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_rcp_f32_e32 v18, v15
	v_div_scale_f32 v21, s0, v10, v8, v10
	v_rcp_f32_e32 v19, v16
	v_div_scale_f32 v20, vcc_lo, v11, v8, v11
	v_div_scale_f32 v22, s1, v9, v8, v9
	v_cvt_f32_i32_e32 v3, v3
	s_delay_alu instid0(TRANS32_DEP_3) | instskip(NEXT) | instid1(TRANS32_DEP_2)
	v_fma_f32 v6, -v14, v17, 1.0
	v_fma_f32 v7, -v15, v18, 1.0
	v_cvt_f32_i32_e32 v0, v0
	v_cvt_f32_i32_e32 v1, v1
	s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f32 v12, -v16, v19, 1.0
	v_dual_fmac_f32 v17, v6, v17 :: v_dual_fmac_f32 v18, v7, v18
	v_add_co_u32 v6, s2, s10, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v19, v12, v19
	v_add_co_ci_u32_e64 v7, s2, s11, v5, s2
	v_dual_mul_f32 v24, v21, v18 :: v_dual_mul_f32 v25, v22, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f32 v13, -v15, v24, v21
	v_mul_f32_e32 v23, v20, v17
	v_fma_f32 v26, -v16, v25, v22
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v24, v13, v18
	v_fma_f32 v12, -v14, v23, v20
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v25, v26, v19
	v_fma_f32 v15, -v15, v24, v21
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v23, v12, v17
	v_fma_f32 v16, -v16, v25, v22
	v_add_co_u32 v12, s2, s12, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e64 v13, s2, s13, v5, s2
	v_fma_f32 v14, -v14, v23, v20
	s_mov_b32 s2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_div_fmas_f32 v14, v14, v17, v23
	s_mov_b32 vcc_lo, s0
	v_div_fmas_f32 v15, v15, v18, v24
	s_mov_b32 vcc_lo, s1
	v_div_fixup_f32 v11, v14, v8, v11
	v_div_fmas_f32 v16, v16, v19, v25
	v_add_co_u32 v4, vcc_lo, s14, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s15, v5, vcc_lo
	v_div_fixup_f32 v10, v15, v8, v10
	s_delay_alu instid0(VALU_DEP_4)
	v_div_fixup_f32 v8, v16, v8, v9
	s_mov_b64 s[0:1], 8
	global_store_b32 v[6:7], v11, off
	global_store_b32 v[12:13], v10, off
	global_store_b32 v[4:5], v8, off
.LBB0_6:                                ; =>This Inner Loop Header: Depth=1
	v_add_nc_u32_e32 v4, s2, v2
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, c@rel32@lo-4
	s_addc_u32 s9, s9, c@rel32@hi+4
	s_add_u32 s8, s0, s8
	s_addc_u32 s9, s1, s9
	s_add_u32 s0, s0, 12
	v_ashrrev_i32_e32 v5, 31, v4
	s_addc_u32 s1, s1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	v_add_co_u32 v4, vcc_lo, s4, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	global_load_b32 v6, v[4:5], off
	s_clause 0x1
	s_load_b64 s[10:11], s[8:9], 0x0
	s_load_b32 s3, s[8:9], 0x8
	s_waitcnt lgkmcnt(0)
	v_dual_add_f32 v5, s11, v0 :: v_dual_add_f32 v4, s10, v3
	v_add_f32_e32 v7, s3, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_i32_f32_e32 v5, v5
	v_cvt_i32_f32_e32 v4, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_i32_f32_e32 v7, v7
	v_add_nc_u32_e32 v5, 0x200, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v4, 0x200, v4
	v_add_nc_u32_e32 v7, 0x200, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v9, 31, v5
	v_ashrrev_i32_e32 v8, 31, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v10, 31, v7
	v_lshrrev_b32_e32 v9, 23, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshrrev_b32_e32 v8, 23, v8
	v_lshrrev_b32_e32 v10, 23, v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v9, v5, v9
	v_add_nc_u32_e32 v8, v4, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v10, v7, v10
	v_and_b32_e32 v9, 0x7ffe00, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v8, 0xfffffe00, v8
	v_and_b32_e32 v10, 0x3e00, v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v5, v5, v9
	v_sub_nc_u32_e32 v4, v4, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_nc_u32_e32 v7, v7, v10
	v_lshlrev_b32_e32 v5, 9, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshl_add_u32 v4, v7, 18, v4
	v_add3_u32 v4, v4, v5, s2
	s_add_i32 s2, s2, 0x8000000
	s_cmpk_lg_i32 s0, 0xec
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v5, 31, v4
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s6, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v5, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[4:5], v6, off
	s_cbranch_scc1 .LBB0_6
.LBB0_7:                                ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10lbm_kernelPfS_S_S_S_S_f
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.amdhsa_user_sgpr_count 13
		.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 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 27
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z10lbm_kernelPfS_S_S_S_S_f, .Lfunc_end0-_Z10lbm_kernelPfS_S_S_S_S_f
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1164
; NumSgprs: 18
; NumVgprs: 27
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 27
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; 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: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z16collision_kernelPfS_S_S_S_S_f ; -- Begin function _Z16collision_kernelPfS_S_S_S_S_f
	.globl	_Z16collision_kernelPfS_S_S_S_S_f
	.p2align	8
	.type	_Z16collision_kernelPfS_S_S_S_S_f,@function
_Z16collision_kernelPfS_S_S_S_S_f:      ; @_Z16collision_kernelPfS_S_S_S_S_f
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x44
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v2, v0, 10, 10
	v_bfe_u32 v3, v0, 20, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[6:7], null, s13, s2, v[1:2]
	v_mad_u64_u32 v[0:1], null, s14, s4, v[2:3]
	v_mad_u64_u32 v[1:2], null, s15, s3, v[3:4]
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max3_i32 v2, v6, v0, v1
	v_cmpx_gt_i32_e32 0x200, v2
	s_cbranch_execz .LBB1_3
; %bb.1:
	v_lshlrev_b32_e32 v7, 18, v1
	v_lshlrev_b32_e32 v10, 9, v0
	s_clause 0x2
	s_load_b128 s[16:19], s[0:1], 0x20
	s_load_b32 s12, s[0:1], 0x30
	s_load_b256 s[4:11], s[0:1], 0x0
	s_mov_b64 s[2:3], 0
	v_add3_u32 v0, v10, v6, v7
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, c@rel32@lo+12
	s_addc_u32 s1, s1, c@rel32@hi+20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
	v_lshlrev_b64 v[8:9], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s16, v8
	v_add_co_ci_u32_e32 v1, vcc_lo, s17, v9, vcc_lo
	v_add_co_u32 v2, vcc_lo, s10, v8
	v_add_co_ci_u32_e32 v3, vcc_lo, s11, v9, vcc_lo
	global_load_b32 v11, v[0:1], off
	v_add_co_u32 v4, vcc_lo, s18, v8
	v_add_co_ci_u32_e32 v5, vcc_lo, s19, v9, vcc_lo
	global_load_b32 v12, v[2:3], off
	global_load_b32 v13, v[4:5], off
	v_add_co_u32 v8, vcc_lo, s8, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s9, v9, vcc_lo
	s_waitcnt vmcnt(2)
	v_mul_f32_e32 v11, v11, v11
	v_add3_u32 v6, v6, v7, v10
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v11, v12, v12
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v11, v13, v13
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f32_e32 v10, 0x3fc00000, v11
.LBB1_2:                                ; =>This Inner Loop Header: Depth=1
	global_load_b32 v15, v[0:1], off
	global_load_b32 v16, v[2:3], off
	global_load_b32 v17, v[4:5], off
	v_ashrrev_i32_e32 v7, 31, v6
	s_add_u32 s8, s0, -8
	s_addc_u32 s9, s1, -1
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, t@rel32@lo+4
	s_addc_u32 s11, s11, t@rel32@hi+12
	s_add_u32 s10, s2, s10
	v_lshlrev_b64 v[11:12], 2, v[6:7]
	global_load_b32 v7, v[8:9], off
	s_addc_u32 s11, s3, s11
	s_add_u32 s2, s2, 4
	s_addc_u32 s3, s3, 0
	v_add_nc_u32_e32 v6, 0x8000000, v6
	v_add_co_u32 v13, vcc_lo, s6, v11
	v_add_co_ci_u32_e32 v14, vcc_lo, s7, v12, vcc_lo
	v_add_co_u32 v11, vcc_lo, s4, v11
	v_add_co_ci_u32_e32 v12, vcc_lo, s5, v12, vcc_lo
	global_load_b32 v13, v[13:14], off
	s_clause 0x1
	s_load_b64 s[8:9], s[8:9], 0x0
	s_load_b32 s13, s[0:1], 0x0
	s_add_u32 s0, s0, 12
	s_addc_u32 s1, s1, 0
	s_cmpk_lg_i32 s2, 0x4c
	s_waitcnt vmcnt(4) lgkmcnt(0)
	v_mul_f32_e32 v14, s9, v15
	s_load_b32 s9, s[10:11], 0x0
	s_waitcnt vmcnt(1) lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v14, s8, v16 :: v_dual_mul_f32 v7, s9, v7
	v_fmac_f32_e32 v14, s13, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f32 v15, 0x40400000, v14, 1.0
	v_mul_f32_e32 v16, 0x40900000, v14
	v_fmac_f32_e32 v15, v14, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v14, v15, v10
	s_waitcnt vmcnt(0)
	v_fma_f32 v7, -v7, v14, v13
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f32 v7, -v7, s12, v13
	global_store_b32 v[11:12], v7, off
	s_cbranch_scc1 .LBB1_2
.LBB1_3:                                ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z16collision_kernelPfS_S_S_S_S_f
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.amdhsa_user_sgpr_count 13
		.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 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 18
		.amdhsa_next_free_sgpr 20
		.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_end1:
	.size	_Z16collision_kernelPfS_S_S_S_S_f, .Lfunc_end1-_Z16collision_kernelPfS_S_S_S_S_f
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 600
; NumSgprs: 22
; NumVgprs: 18
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 18
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; 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: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	c                       ; @c
	.type	c,@object
	.section	.bss,"aw",@nobits
	.globl	c
	.p2align	4, 0x0
c:
	.zero	228
	.size	c, 228

	.protected	t                       ; @t
	.type	t,@object
	.globl	t
	.p2align	4, 0x0
t:
	.zero	76
	.size	t, 76

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 c
	.addrsig_sym t
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         60
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         68
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         70
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         72
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         74
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         76
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         78
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         120
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 312
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10lbm_kernelPfS_S_S_S_S_f
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10lbm_kernelPfS_S_S_S_S_f.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     27
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         60
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         68
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         70
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         72
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         74
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         76
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         78
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         120
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 312
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z16collision_kernelPfS_S_S_S_S_f
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z16collision_kernelPfS_S_S_S_S_f.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     18
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
