// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args "-Xclang" "-fcuda-allow-variadic-functions"

/*
Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
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.
*/

// CHECK: #include "hip/hip_runtime.h"
#include "cuda_runtime.h"

#include <algorithm>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <type_traits>

#define EXIT_WAIVED 2

const char* sampleName = "hipSimpleAtomicsTest";

using namespace std;
// Auto-Verification Code
bool testResult = true;

bool computeGoldBitwise(...) {
    return true;
}

template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
bool computeGoldBitwise(T* gpuData, int len) {
    T val = 0xff;
    for (int i = 0; i < len; ++i) {
        // 9th element should be 1
        val &= (2 * i + 7);
    }
    if (val != gpuData[8]) {
        printf("atomicAnd failed\n");
        return false;
    }

    val = 0;
    for (int i = 0; i < len; ++i) {
        // 10th element should be 0xff
        val |= (1 << i);
    }
    if (val != gpuData[9]) {
        printf("atomicOr failed\n");
        return false;
    }

    val = 0xff;
    for (int i = 0; i < len; ++i) {
        // 11th element should be 0xff
        val ^= i;
    }
    if (val != gpuData[10]) {
        printf("atomicXor failed\n");
        return false;
    }

    return true;
}

template<typename T>
bool computeGold(T* gpuData, int len) {
    T val = 0;
    for (int i = 0; i < len; ++i) {
        val += 10;
    }
    if (val != gpuData[0]) {
        printf("atomicAdd failed\n");
        return false;
    }

    val = 0;
    for (int i = 0; i < len; ++i) {
        val -= 10;
    }
    if (val != gpuData[1]) {
        printf("atomicSub failed\n");
        return false;
    }

    bool found = false;
    for (T i = 0; i < len; ++i) {
        // third element should be a member of [0, len)
        if (i == gpuData[2]) {
            found = true;
            break;
        }
    }
    if (!found) {
        printf("atomicExch failed\n");
        return false;
    }

    val = -(1 << 8);
    for (T i = 0; i < len; ++i) {
        // fourth element should be len-1
        val = max(val, i);
    }
    if (val != gpuData[3]) {
        printf("atomicMax failed\n");
        return false;
    }

    val = 1 << 8;
    for (T i = 0; i < len; ++i) {
        val = min(val, i);
    }
    if (val != gpuData[4]) {
        printf("atomicMin failed\n");
        return false;
    }

    int limit = 17;
    val = 0;
    for (int i = 0; i < len; ++i) {
        val = (val >= limit) ? 0 : val + 1;
    }
    if (val != gpuData[5]) {
        printf("atomicInc failed\n");
        return false;
    }

    limit = 137;
    val = 0;
    for (int i = 0; i < len; ++i) {
        val = ((val == 0) || (val > limit)) ? limit : val - 1;
    }
    if (val != gpuData[6]) {
        printf("atomicDec failed\n");
        return false;
    }

    found = false;
    for (T i = 0; i < len; ++i) {
        // eighth element should be a member of [0, len)
        if (i == gpuData[7]) {
            found = true;
            break;
        }
    }
    if (!found) {
        printf("atomicCAS failed\n");
        return false;
    }

    return computeGoldBitwise(gpuData, len);
}

__device__
void testKernelExch(...) {}

template<typename T, typename enable_if<!is_same<T, double>{}>::type* = nullptr>
__device__
void testKernelExch(T* g_odata) {
    // access thread id
    const T tid = blockDim.x * blockIdx.x + threadIdx.x;
    // Atomic exchange
    atomicExch(&g_odata[2], tid);
}

__device__
void testKernelSub(...) {}

template<
    typename T,
    typename enable_if<
        is_same<T, int>{} || is_same<T, unsigned int>{}>::type* = nullptr>
__device__
void testKernelSub(T* g_odata) {
    // Atomic subtraction (final should be 0)
    atomicSub(&g_odata[1], 10);
}

__device__
void testKernelIntegral(...) {}

template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
__device__
void testKernelIntegral(T* g_odata) {
    // access thread id
    const T tid = blockDim.x * blockIdx.x + threadIdx.x;
    // Atomic maximum
    atomicMax(&g_odata[3], tid);
    // Atomic minimum
    atomicMin(&g_odata[4], tid);
    // Atomic increment (modulo 17+1)
    atomicInc((unsigned int*)&g_odata[5], 17);
    // Atomic decrement
    atomicDec((unsigned int*)&g_odata[6], 137);
    // Atomic compare-and-swap
    atomicCAS(&g_odata[7], tid - 1, tid);
    // Atomic AND
    atomicAnd(&g_odata[8], 2 * tid + 7);
    // Atomic OR
    atomicOr(&g_odata[9], 1 << tid);
    // Atomic XOR
    atomicXor(&g_odata[10], tid);
    testKernelSub(g_odata);
}

template<typename T>
__global__ void testKernel(T* g_odata) {
    // Atomic addition
    atomicAdd(&g_odata[0], 10);
    testKernelIntegral(g_odata);
    testKernelExch(g_odata);
}

template<typename T>
void runTest() {
    // CHECK: hipDeviceProp_t deviceProp;
    cudaDeviceProp deviceProp;
    deviceProp.major = 0;
    deviceProp.minor = 0;
    int dev = 0;
    // CHECK: hipGetDeviceProperties(&deviceProp, dev);
    cudaGetDeviceProperties(&deviceProp, dev);
    // Statistics about the GPU device
    printf(
        "> GPU device has %d Multi-Processors, "
        "SM %d.%d compute capabilities\n\n",
        deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
    unsigned int numThreads = 256;
    unsigned int numBlocks = 64;
    unsigned int numData = 11;
    unsigned int memSize = sizeof(T) * numData;

    // Allocate mem for the result on host side
    T* hOData = (T*)malloc(memSize);
    // Initialize the memory
    for (unsigned int i = 0; i < numData; i++) hOData[i] = 0;
    // To make the AND and XOR tests generate something other than 0...
    hOData[8] = hOData[10] = 0xff;
    // Allocate device memory for result
    T* dOData;
    // CHECK: hipMalloc((void**)&dOData, memSize);
    cudaMalloc((void**)&dOData, memSize);
    // Copy host memory to device to initialize to zero
    // CHECK: hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice);
    cudaMemcpy(dOData, hOData, memSize, cudaMemcpyHostToDevice);
    // Execute the kernel
    // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(testKernel), dim3(numBlocks), dim3(numThreads), 0, 0, dOData);
    testKernel<<<numBlocks, numThreads>>>(dOData);
    // Copy result from device to host
    // CHECK: hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost);
    cudaMemcpy(hOData, dOData, memSize, cudaMemcpyDeviceToHost);
    // Compute reference solution
    testResult = computeGold(hOData, numThreads * numBlocks);
    // Cleanup memory
    free(hOData);
    // CHECK: hipFree(dOData);
    cudaFree(dOData);
}

int main(int argc, char** argv) {
    printf("%s starting...\n", sampleName);
    runTest<int>();
    runTest<unsigned int>();
    runTest<unsigned long long>();
    runTest<float>();
    runTest<double>();
    // CHECK: hipDeviceReset();
    cudaDeviceReset();
    printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!");
    exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
}

	.file	"tmpxft_002a3abd_00000000-6_atomics.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2322:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2322:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZL35__device_stub__Z10testKernelIdEvPT_Pd, @function
_ZL35__device_stub__Z10testKernelIdEvPT_Pd:
.LFB2352:
	.cfi_startproc
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z10testKernelIdEvPT_(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L2:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2352:
	.size	_ZL35__device_stub__Z10testKernelIdEvPT_Pd, .-_ZL35__device_stub__Z10testKernelIdEvPT_Pd
	.section	.text._Z10testKernelIdEvPT_,"axG",@progbits,_Z10testKernelIdEvPT_,comdat
	.weak	_Z10testKernelIdEvPT_
	.type	_Z10testKernelIdEvPT_, @function
_Z10testKernelIdEvPT_:
.LFB2411:
	.cfi_startproc
	endbr64
	jmp	_ZL35__device_stub__Z10testKernelIdEvPT_Pd
	.cfi_endproc
.LFE2411:
	.size	_Z10testKernelIdEvPT_, .-_Z10testKernelIdEvPT_
	.text
	.type	_ZL35__device_stub__Z10testKernelIiEvPT_Pi, @function
_ZL35__device_stub__Z10testKernelIiEvPT_Pi:
.LFB2344:
	.cfi_startproc
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z10testKernelIiEvPT_(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L8:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2344:
	.size	_ZL35__device_stub__Z10testKernelIiEvPT_Pi, .-_ZL35__device_stub__Z10testKernelIiEvPT_Pi
	.section	.text._Z10testKernelIiEvPT_,"axG",@progbits,_Z10testKernelIiEvPT_,comdat
	.weak	_Z10testKernelIiEvPT_
	.type	_Z10testKernelIiEvPT_, @function
_Z10testKernelIiEvPT_:
.LFB2406:
	.cfi_startproc
	endbr64
	jmp	_ZL35__device_stub__Z10testKernelIiEvPT_Pi
	.cfi_endproc
.LFE2406:
	.size	_Z10testKernelIiEvPT_, .-_Z10testKernelIiEvPT_
	.text
	.type	_ZL35__device_stub__Z10testKernelIjEvPT_Pj, @function
_ZL35__device_stub__Z10testKernelIjEvPT_Pj:
.LFB2346:
	.cfi_startproc
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z10testKernelIjEvPT_(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L13:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2346:
	.size	_ZL35__device_stub__Z10testKernelIjEvPT_Pj, .-_ZL35__device_stub__Z10testKernelIjEvPT_Pj
	.section	.text._Z10testKernelIjEvPT_,"axG",@progbits,_Z10testKernelIjEvPT_,comdat
	.weak	_Z10testKernelIjEvPT_
	.type	_Z10testKernelIjEvPT_, @function
_Z10testKernelIjEvPT_:
.LFB2408:
	.cfi_startproc
	endbr64
	jmp	_ZL35__device_stub__Z10testKernelIjEvPT_Pj
	.cfi_endproc
.LFE2408:
	.size	_Z10testKernelIjEvPT_, .-_Z10testKernelIjEvPT_
	.text
	.type	_ZL35__device_stub__Z10testKernelIyEvPT_Py, @function
_ZL35__device_stub__Z10testKernelIyEvPT_Py:
.LFB2348:
	.cfi_startproc
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z10testKernelIyEvPT_(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L18:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2348:
	.size	_ZL35__device_stub__Z10testKernelIyEvPT_Py, .-_ZL35__device_stub__Z10testKernelIyEvPT_Py
	.section	.text._Z10testKernelIyEvPT_,"axG",@progbits,_Z10testKernelIyEvPT_,comdat
	.weak	_Z10testKernelIyEvPT_
	.type	_Z10testKernelIyEvPT_, @function
_Z10testKernelIyEvPT_:
.LFB2409:
	.cfi_startproc
	endbr64
	jmp	_ZL35__device_stub__Z10testKernelIyEvPT_Py
	.cfi_endproc
.LFE2409:
	.size	_Z10testKernelIyEvPT_, .-_Z10testKernelIyEvPT_
	.text
	.type	_ZL35__device_stub__Z10testKernelIfEvPT_Pf, @function
_ZL35__device_stub__Z10testKernelIfEvPT_Pf:
.LFB2350:
	.cfi_startproc
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	leaq	44(%rsp), %rsi
	leaq	32(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 80(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	24(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z10testKernelIfEvPT_(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 128
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	96(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L23:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2350:
	.size	_ZL35__device_stub__Z10testKernelIfEvPT_Pf, .-_ZL35__device_stub__Z10testKernelIfEvPT_Pf
	.section	.text._Z10testKernelIfEvPT_,"axG",@progbits,_Z10testKernelIfEvPT_,comdat
	.weak	_Z10testKernelIfEvPT_
	.type	_Z10testKernelIfEvPT_, @function
_Z10testKernelIfEvPT_:
.LFB2410:
	.cfi_startproc
	endbr64
	jmp	_ZL35__device_stub__Z10testKernelIfEvPT_Pf
	.cfi_endproc
.LFE2410:
	.size	_Z10testKernelIfEvPT_, .-_Z10testKernelIfEvPT_
	.text
	.globl	_Z18computeGoldBitwisez
	.type	_Z18computeGoldBitwisez, @function
_Z18computeGoldBitwisez:
.LFB2307:
	.cfi_startproc
	endbr64
	movb	$1, %al
	ret
	.cfi_endproc
.LFE2307:
	.size	_Z18computeGoldBitwisez, .-_Z18computeGoldBitwisez
	.globl	_Z14testKernelExchz
	.type	_Z14testKernelExchz, @function
_Z14testKernelExchz:
.LFB2310:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2310:
	.size	_Z14testKernelExchz, .-_Z14testKernelExchz
	.globl	_Z13testKernelSubz
	.type	_Z13testKernelSubz, @function
_Z13testKernelSubz:
.LFB2492:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	call	_Z14testKernelExchz
	.cfi_endproc
.LFE2492:
	.size	_Z13testKernelSubz, .-_Z13testKernelSubz
	.globl	_Z18testKernelIntegralz
	.type	_Z18testKernelIntegralz, @function
_Z18testKernelIntegralz:
.LFB2494:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	call	_Z14testKernelExchz
	.cfi_endproc
.LFE2494:
	.size	_Z18testKernelIntegralz, .-_Z18testKernelIntegralz
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z10testKernelIdEvPT_"
.LC1:
	.string	"_Z10testKernelIfEvPT_"
.LC2:
	.string	"_Z10testKernelIyEvPT_"
.LC3:
	.string	"_Z10testKernelIjEvPT_"
.LC4:
	.string	"_Z10testKernelIiEvPT_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2355:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10testKernelIdEvPT_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10testKernelIfEvPT_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10testKernelIyEvPT_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10testKernelIjEvPT_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10testKernelIiEvPT_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2355:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._Z11computeGoldIiEbPT_i.str1.1,"aMS",@progbits,1
.LC5:
	.string	"atomicAdd failed\n"
.LC6:
	.string	"atomicSub failed\n"
.LC7:
	.string	"atomicMax failed\n"
.LC8:
	.string	"atomicMin failed\n"
.LC9:
	.string	"atomicInc failed\n"
.LC10:
	.string	"atomicDec failed\n"
.LC11:
	.string	"atomicAnd failed\n"
.LC12:
	.string	"atomicOr failed\n"
.LC13:
	.string	"atomicXor failed\n"
.LC14:
	.string	"atomicCAS failed\n"
.LC15:
	.string	"atomicExch failed\n"
	.section	.text._Z11computeGoldIiEbPT_i,"axG",@progbits,_Z11computeGoldIiEbPT_i,comdat
	.weak	_Z11computeGoldIiEbPT_i
	.type	_Z11computeGoldIiEbPT_i, @function
_Z11computeGoldIiEbPT_i:
.LFB2450:
	.cfi_startproc
	endbr64
	pushq	%rcx
	.cfi_def_cfa_offset 16
	xorl	%ecx, %ecx
	testl	%esi, %esi
	cmovns	%esi, %ecx
	imull	$10, %ecx, %ecx
	cmpl	%ecx, (%rdi)
	je	.L38
	leaq	.LC5(%rip), %rsi
	jmp	.L80
.L38:
	movl	%esi, %edx
	xorl	%esi, %esi
	addl	4(%rdi), %ecx
	movq	%rdi, %rax
	je	.L40
	leaq	.LC6(%rip), %rsi
.L80:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edx, %edx
	jmp	.L37
.L43:
	cmpl	%esi, 8(%rax)
	je	.L64
	incl	%esi
.L40:
	cmpl	%edx, %esi
	jl	.L43
	jmp	.L81
.L64:
	xorl	%ecx, %ecx
	movl	$-256, %esi
.L42:
	cmpl	%edx, %ecx
	jge	.L82
	cmpl	%ecx, %esi
	cmovl	%ecx, %esi
	incl	%ecx
	jmp	.L42
.L82:
	cmpl	%esi, 12(%rax)
	je	.L65
	leaq	.LC7(%rip), %rsi
	jmp	.L80
.L65:
	xorl	%ecx, %ecx
	movl	$256, %esi
.L46:
	cmpl	%edx, %ecx
	jge	.L83
	cmpl	%ecx, %esi
	cmovg	%ecx, %esi
	incl	%ecx
	jmp	.L46
.L83:
	cmpl	%esi, 16(%rax)
	je	.L66
	leaq	.LC8(%rip), %rsi
	jmp	.L80
.L66:
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L48:
	cmpl	%edx, %esi
	jge	.L84
	cmpl	$17, %ecx
	je	.L67
	incl	%ecx
	jmp	.L49
.L67:
	xorl	%ecx, %ecx
.L49:
	incl	%esi
	jmp	.L48
.L84:
	leaq	.LC9(%rip), %rsi
	cmpl	%ecx, 20(%rax)
	jne	.L80
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L51:
	cmpl	%edx, %esi
	jge	.L85
	testl	%ecx, %ecx
	je	.L69
	decl	%ecx
	jmp	.L52
.L69:
	movl	$137, %ecx
.L52:
	incl	%esi
	jmp	.L51
.L85:
	xorl	%esi, %esi
	cmpl	%ecx, 24(%rax)
	je	.L54
	leaq	.LC10(%rip), %rsi
	jmp	.L80
.L56:
	cmpl	%esi, 28(%rax)
	je	.L71
	incl	%esi
.L54:
	cmpl	%edx, %esi
	jl	.L56
	jmp	.L86
.L71:
	xorl	%ecx, %ecx
	movl	$255, %esi
.L55:
	cmpl	%ecx, %edx
	jle	.L87
	leal	7(%rcx,%rcx), %edi
	incl	%ecx
	andl	%edi, %esi
	jmp	.L55
.L87:
	cmpl	32(%rax), %esi
	je	.L72
	leaq	.LC11(%rip), %rsi
	jmp	.L80
.L72:
	xorl	%ecx, %ecx
	xorl	%esi, %esi
	movl	$1, %edi
.L59:
	cmpl	%ecx, %edx
	jle	.L88
	movl	%edi, %r8d
	sall	%cl, %r8d
	incl	%ecx
	orl	%r8d, %esi
	jmp	.L59
.L88:
	cmpl	36(%rax), %esi
	je	.L73
	leaq	.LC12(%rip), %rsi
	jmp	.L80
.L73:
	xorl	%ecx, %ecx
	movl	$255, %esi
.L61:
	cmpl	%ecx, %edx
	jle	.L89
	xorl	%ecx, %esi
	incl	%ecx
	jmp	.L61
.L89:
	movb	$1, %dl
	cmpl	40(%rax), %esi
	je	.L37
	leaq	.LC13(%rip), %rsi
	jmp	.L80
.L86:
	leaq	.LC14(%rip), %rsi
	jmp	.L80
.L81:
	leaq	.LC15(%rip), %rsi
	jmp	.L80
.L37:
	movl	%edx, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2450:
	.size	_Z11computeGoldIiEbPT_i, .-_Z11computeGoldIiEbPT_i
	.section	.rodata._Z7runTestIiEvv.str1.1,"aMS",@progbits,1
.LC16:
	.string	"> GPU device has %d Multi-Processors, SM %d.%d compute capabilities\n\n"
	.section	.text._Z7runTestIiEvv,"axG",@progbits,_Z7runTestIiEvv,comdat
	.weak	_Z7runTestIiEvv
	.type	_Z7runTestIiEvv, @function
_Z7runTestIiEvv:
.LFB2401:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%esi, %esi
	subq	$1072, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1064(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rax, 392(%rsp)
	call	cudaGetDeviceProperties_v2@PLT
	movl	396(%rsp), %r8d
	xorl	%eax, %eax
	movl	420(%rsp), %edx
	movl	392(%rsp), %ecx
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$44, %edi
	call	malloc@PLT
	movl	$11, %ecx
	movl	$44, %esi
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	movl	$255, 40(%rbx)
	movq	%rsp, %rdi
	movl	$255, 32(%rbx)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$44, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$67108865, %edi
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$6, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L91
	movq	(%rsp), %rdi
	call	_ZL35__device_stub__Z10testKernelIiEvPT_Pi
.L91:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$44, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$16384, %esi
	movq	%rbx, %rdi
	call	_Z11computeGoldIiEbPT_i
	movq	%rbx, %rdi
	movb	%al, testResult(%rip)
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	1064(%rsp), %rax
	subq	%fs:40, %rax
	je	.L92
	call	__stack_chk_fail@PLT
.L92:
	addq	$1072, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2401:
	.size	_Z7runTestIiEvv, .-_Z7runTestIiEvv
	.section	.text._Z11computeGoldIjEbPT_i,"axG",@progbits,_Z11computeGoldIjEbPT_i,comdat
	.weak	_Z11computeGoldIjEbPT_i
	.type	_Z11computeGoldIjEbPT_i, @function
_Z11computeGoldIjEbPT_i:
.LFB2452:
	.cfi_startproc
	endbr64
	pushq	%rcx
	.cfi_def_cfa_offset 16
	xorl	%ecx, %ecx
	testl	%esi, %esi
	cmovns	%esi, %ecx
	imull	$10, %ecx, %ecx
	cmpl	%ecx, (%rdi)
	je	.L95
	leaq	.LC5(%rip), %rsi
	jmp	.L134
.L95:
	addl	4(%rdi), %ecx
	movq	%rdi, %rax
	jne	.L97
	movl	%esi, %edx
	xorl	%ecx, %ecx
	jmp	.L98
.L97:
	leaq	.LC6(%rip), %rsi
.L134:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edx, %edx
	jmp	.L94
.L101:
	cmpl	%ecx, 8(%rax)
	je	.L120
	incl	%ecx
.L98:
	cmpl	%ecx, %edx
	jne	.L101
	jmp	.L135
.L120:
	xorl	%ecx, %ecx
	movl	$-256, %esi
.L100:
	cmpl	%ecx, %edx
	je	.L136
	cmpl	%ecx, %esi
	cmovb	%ecx, %esi
	incl	%ecx
	jmp	.L100
.L136:
	cmpl	%esi, 12(%rax)
	je	.L121
	leaq	.LC7(%rip), %rsi
	jmp	.L134
.L121:
	xorl	%ecx, %ecx
	movl	$256, %esi
.L104:
	cmpl	%ecx, %edx
	je	.L137
	cmpl	%ecx, %esi
	cmova	%ecx, %esi
	incl	%ecx
	jmp	.L104
.L137:
	cmpl	%esi, 16(%rax)
	je	.L122
	leaq	.LC8(%rip), %rsi
	jmp	.L134
.L122:
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L106:
	cmpl	%edx, %esi
	jge	.L138
	cmpl	$17, %ecx
	je	.L123
	incl	%ecx
	jmp	.L107
.L123:
	xorl	%ecx, %ecx
.L107:
	incl	%esi
	jmp	.L106
.L138:
	leaq	.LC9(%rip), %rsi
	cmpl	%ecx, 20(%rax)
	jne	.L134
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L109:
	cmpl	%edx, %esi
	jge	.L139
	decl	%ecx
	movl	$137, %edi
	cmpl	%edi, %ecx
	cmova	%edi, %ecx
	incl	%esi
	jmp	.L109
.L139:
	xorl	%esi, %esi
	cmpl	%ecx, 24(%rax)
	je	.L111
	leaq	.LC10(%rip), %rsi
	jmp	.L134
.L113:
	cmpl	%esi, 28(%rax)
	je	.L126
	incl	%esi
.L111:
	cmpl	%esi, %edx
	jne	.L113
	jmp	.L140
.L126:
	xorl	%ecx, %ecx
	movl	$255, %esi
.L112:
	cmpl	%ecx, %edx
	jle	.L141
	leal	7(%rcx,%rcx), %edi
	incl	%ecx
	andl	%edi, %esi
	jmp	.L112
.L141:
	cmpl	32(%rax), %esi
	je	.L127
	leaq	.LC11(%rip), %rsi
	jmp	.L134
.L127:
	xorl	%ecx, %ecx
	xorl	%esi, %esi
	movl	$1, %edi
.L116:
	cmpl	%ecx, %edx
	jle	.L142
	movl	%edi, %r8d
	sall	%cl, %r8d
	incl	%ecx
	orl	%r8d, %esi
	jmp	.L116
.L142:
	cmpl	36(%rax), %esi
	je	.L128
	leaq	.LC12(%rip), %rsi
	jmp	.L134
.L128:
	xorl	%ecx, %ecx
	movl	$255, %esi
.L118:
	cmpl	%ecx, %edx
	jle	.L143
	xorl	%ecx, %esi
	incl	%ecx
	jmp	.L118
.L143:
	movb	$1, %dl
	cmpl	40(%rax), %esi
	je	.L94
	leaq	.LC13(%rip), %rsi
	jmp	.L134
.L140:
	leaq	.LC14(%rip), %rsi
	jmp	.L134
.L135:
	leaq	.LC15(%rip), %rsi
	jmp	.L134
.L94:
	movl	%edx, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2452:
	.size	_Z11computeGoldIjEbPT_i, .-_Z11computeGoldIjEbPT_i
	.section	.text._Z7runTestIjEvv,"axG",@progbits,_Z7runTestIjEvv,comdat
	.weak	_Z7runTestIjEvv
	.type	_Z7runTestIjEvv, @function
_Z7runTestIjEvv:
.LFB2402:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%esi, %esi
	subq	$1072, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1064(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rax, 392(%rsp)
	call	cudaGetDeviceProperties_v2@PLT
	movl	396(%rsp), %r8d
	xorl	%eax, %eax
	movl	420(%rsp), %edx
	movl	392(%rsp), %ecx
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$44, %edi
	call	malloc@PLT
	movl	$11, %ecx
	movl	$44, %esi
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	movl	$255, 40(%rbx)
	movq	%rsp, %rdi
	movl	$255, 32(%rbx)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$44, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$67108865, %edi
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$6, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L145
	movq	(%rsp), %rdi
	call	_ZL35__device_stub__Z10testKernelIjEvPT_Pj
.L145:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$44, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$16384, %esi
	movq	%rbx, %rdi
	call	_Z11computeGoldIjEbPT_i
	movq	%rbx, %rdi
	movb	%al, testResult(%rip)
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	1064(%rsp), %rax
	subq	%fs:40, %rax
	je	.L146
	call	__stack_chk_fail@PLT
.L146:
	addq	$1072, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2402:
	.size	_Z7runTestIjEvv, .-_Z7runTestIjEvv
	.section	.text._Z11computeGoldIyEbPT_i,"axG",@progbits,_Z11computeGoldIyEbPT_i,comdat
	.weak	_Z11computeGoldIyEbPT_i
	.type	_Z11computeGoldIyEbPT_i, @function
_Z11computeGoldIyEbPT_i:
.LFB2453:
	.cfi_startproc
	endbr64
	movslq	%esi, %rdx
	pushq	%rcx
	.cfi_def_cfa_offset 16
	xorl	%ecx, %ecx
	leaq	.LC5(%rip), %rsi
	testl	%edx, %edx
	cmovns	%rdx, %rcx
	imulq	$10, %rcx, %rcx
	cmpq	%rcx, (%rdi)
	jne	.L191
	movq	%rdi, %rax
	xorl	%esi, %esi
	movslq	%edx, %rdi
	addq	8(%rax), %rcx
	je	.L151
	leaq	.LC6(%rip), %rsi
.L191:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edx, %edx
	jmp	.L148
.L154:
	cmpq	%rsi, 16(%rax)
	je	.L175
	incq	%rsi
.L151:
	cmpq	%rdi, %rsi
	jb	.L154
	jmp	.L192
.L175:
	xorl	%ecx, %ecx
	movq	$-256, %rsi
.L153:
	cmpq	%rcx, %rdi
	je	.L193
	cmpq	%rcx, %rsi
	cmovb	%rcx, %rsi
	incq	%rcx
	jmp	.L153
.L193:
	cmpq	%rsi, 24(%rax)
	je	.L176
	leaq	.LC7(%rip), %rsi
	jmp	.L191
.L176:
	xorl	%ecx, %ecx
	movl	$256, %esi
.L157:
	cmpq	%rcx, %rdi
	je	.L194
	cmpq	%rcx, %rsi
	cmova	%rcx, %rsi
	incq	%rcx
	jmp	.L157
.L194:
	cmpq	%rsi, 32(%rax)
	je	.L177
	leaq	.LC8(%rip), %rsi
	jmp	.L191
.L177:
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L159:
	cmpl	%edx, %esi
	jge	.L195
	cmpq	$17, %rcx
	je	.L178
	incq	%rcx
	jmp	.L160
.L178:
	xorl	%ecx, %ecx
.L160:
	incl	%esi
	jmp	.L159
.L195:
	leaq	.LC9(%rip), %rsi
	cmpq	%rcx, 40(%rax)
	jne	.L191
	xorl	%esi, %esi
	xorl	%ecx, %ecx
.L162:
	cmpl	%edx, %esi
	jge	.L196
	testq	%rcx, %rcx
	je	.L180
	decq	%rcx
	jmp	.L163
.L180:
	movl	$137, %ecx
.L163:
	incl	%esi
	jmp	.L162
.L196:
	xorl	%esi, %esi
	cmpq	%rcx, 48(%rax)
	je	.L165
	leaq	.LC10(%rip), %rsi
	jmp	.L191
.L167:
	cmpq	%rsi, 56(%rax)
	je	.L182
	incq	%rsi
.L165:
	cmpq	%rsi, %rdi
	jne	.L167
	jmp	.L197
.L182:
	xorl	%ecx, %ecx
	movl	$255, %esi
.L166:
	cmpl	%ecx, %edx
	jle	.L198
	leaq	(%rcx,%rcx), %rdi
	incq	%rcx
	addq	$7, %rdi
	andq	%rdi, %rsi
	jmp	.L166
.L198:
	cmpq	64(%rax), %rsi
	je	.L183
	leaq	.LC11(%rip), %rsi
	jmp	.L191
.L183:
	xorl	%ecx, %ecx
	xorl	%esi, %esi
	movl	$1, %r8d
.L170:
	cmpl	%ecx, %edx
	jle	.L199
	movl	%r8d, %edi
	sall	%cl, %edi
	incl	%ecx
	movslq	%edi, %rdi
	orq	%rdi, %rsi
	jmp	.L170
.L199:
	cmpq	72(%rax), %rsi
	je	.L184
	leaq	.LC12(%rip), %rsi
	jmp	.L191
.L184:
	xorl	%ecx, %ecx
	movl	$255, %esi
.L172:
	cmpl	%ecx, %edx
	jle	.L200
	xorq	%rcx, %rsi
	incq	%rcx
	jmp	.L172
.L200:
	movb	$1, %dl
	cmpq	80(%rax), %rsi
	je	.L148
	leaq	.LC13(%rip), %rsi
	jmp	.L191
.L197:
	leaq	.LC14(%rip), %rsi
	jmp	.L191
.L192:
	leaq	.LC15(%rip), %rsi
	jmp	.L191
.L148:
	movl	%edx, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2453:
	.size	_Z11computeGoldIyEbPT_i, .-_Z11computeGoldIyEbPT_i
	.section	.text._Z7runTestIyEvv,"axG",@progbits,_Z7runTestIyEvv,comdat
	.weak	_Z7runTestIyEvv
	.type	_Z7runTestIyEvv, @function
_Z7runTestIyEvv:
.LFB2403:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%esi, %esi
	subq	$1072, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1064(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rax, 392(%rsp)
	call	cudaGetDeviceProperties_v2@PLT
	movl	396(%rsp), %r8d
	xorl	%eax, %eax
	movl	420(%rsp), %edx
	movl	392(%rsp), %ecx
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$88, %edi
	call	malloc@PLT
	movl	$22, %ecx
	movl	$88, %esi
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	movq	$255, 80(%rbx)
	movq	%rsp, %rdi
	movq	$255, 64(%rbx)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$88, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$67108865, %edi
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$6, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L202
	movq	(%rsp), %rdi
	call	_ZL35__device_stub__Z10testKernelIyEvPT_Py
.L202:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$88, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$16384, %esi
	movq	%rbx, %rdi
	call	_Z11computeGoldIyEbPT_i
	movq	%rbx, %rdi
	movb	%al, testResult(%rip)
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	1064(%rsp), %rax
	subq	%fs:40, %rax
	je	.L203
	call	__stack_chk_fail@PLT
.L203:
	addq	$1072, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2403:
	.size	_Z7runTestIyEvv, .-_Z7runTestIyEvv
	.section	.text._Z11computeGoldIfEbPT_i,"axG",@progbits,_Z11computeGoldIfEbPT_i,comdat
	.weak	_Z11computeGoldIfEbPT_i
	.type	_Z11computeGoldIfEbPT_i, @function
_Z11computeGoldIfEbPT_i:
.LFB2454:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%eax, %eax
	movl	%esi, %ebp
	xorps	%xmm0, %xmm0
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movss	.LC21(%rip), %xmm1
.L206:
	cmpl	%ebp, %eax
	jge	.L262
	addss	%xmm1, %xmm0
	incl	%eax
	jmp	.L206
.L262:
	ucomiss	(%rbx), %xmm0
	jp	.L247
	je	.L237
.L247:
	leaq	.LC5(%rip), %rsi
	jmp	.L261
.L237:
	movss	.LC21(%rip), %xmm1
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L208:
	cmpl	%ebp, %eax
	jge	.L263
	subss	%xmm1, %xmm0
	incl	%eax
	jmp	.L208
.L263:
	ucomiss	4(%rbx), %xmm0
	jp	.L248
	xorps	%xmm0, %xmm0
	jne	.L248
	cvtsi2ssl	%ebp, %xmm2
	movss	.LC22(%rip), %xmm1
	jmp	.L212
.L248:
	leaq	.LC6(%rip), %rsi
.L261:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%eax, %eax
	jmp	.L205
.L217:
	ucomiss	8(%rbx), %xmm0
	jnp	.L264
.L249:
	addss	%xmm1, %xmm0
.L212:
	comiss	%xmm0, %xmm2
	ja	.L217
	jmp	.L265
.L264:
	jne	.L249
	movss	.LC18(%rip), %xmm0
	xorps	%xmm1, %xmm1
.L215:
	comiss	%xmm1, %xmm2
	jbe	.L266
	movss	%xmm2, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	call	fmaxf@PLT
	movss	8(%rsp), %xmm1
	movss	12(%rsp), %xmm2
	addss	.LC22(%rip), %xmm1
	jmp	.L215
.L266:
	ucomiss	12(%rbx), %xmm0
	jp	.L250
	je	.L240
.L250:
	leaq	.LC7(%rip), %rsi
	jmp	.L261
.L240:
	movss	.LC19(%rip), %xmm0
	xorps	%xmm1, %xmm1
.L220:
	comiss	%xmm1, %xmm2
	jbe	.L267
	movss	%xmm2, 12(%rsp)
	movss	%xmm1, 8(%rsp)
	call	fminf@PLT
	movss	8(%rsp), %xmm1
	movss	12(%rsp), %xmm2
	addss	.LC22(%rip), %xmm1
	jmp	.L220
.L267:
	ucomiss	16(%rbx), %xmm0
	jp	.L251
	je	.L241
.L251:
	leaq	.LC8(%rip), %rsi
	jmp	.L261
.L241:
	movss	.LC22(%rip), %xmm1
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L223:
	cmpl	%ebp, %eax
	jge	.L268
	comiss	.LC23(%rip), %xmm0
	jnb	.L242
	addss	%xmm1, %xmm0
	jmp	.L225
.L242:
	xorps	%xmm0, %xmm0
.L225:
	incl	%eax
	jmp	.L223
.L268:
	ucomiss	20(%rbx), %xmm0
	jp	.L252
	je	.L243
.L252:
	leaq	.LC9(%rip), %rsi
	jmp	.L261
.L243:
	xorps	%xmm0, %xmm0
	movss	.LC22(%rip), %xmm3
	xorl	%eax, %eax
	movaps	%xmm0, %xmm1
.L227:
	cmpl	%ebp, %eax
	jge	.L269
	ucomiss	%xmm1, %xmm0
	jp	.L253
	je	.L244
.L253:
	subss	%xmm3, %xmm0
	jmp	.L229
.L244:
	movss	.LC20(%rip), %xmm0
.L229:
	incl	%eax
	jmp	.L227
.L269:
	ucomiss	24(%rbx), %xmm0
	jp	.L254
	movss	.LC22(%rip), %xmm1
	xorps	%xmm0, %xmm0
	je	.L232
.L254:
	leaq	.LC10(%rip), %rsi
	jmp	.L261
.L235:
	ucomiss	28(%rbx), %xmm0
	jnp	.L270
.L255:
	addss	%xmm1, %xmm0
.L232:
	comiss	%xmm0, %xmm2
	ja	.L235
	jmp	.L271
.L270:
	jne	.L255
	movb	$1, %al
	jmp	.L205
.L271:
	leaq	.LC14(%rip), %rsi
	jmp	.L261
.L265:
	leaq	.LC15(%rip), %rsi
	jmp	.L261
.L205:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2454:
	.size	_Z11computeGoldIfEbPT_i, .-_Z11computeGoldIfEbPT_i
	.section	.text._Z7runTestIfEvv,"axG",@progbits,_Z7runTestIfEvv,comdat
	.weak	_Z7runTestIfEvv
	.type	_Z7runTestIfEvv, @function
_Z7runTestIfEvv:
.LFB2404:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%esi, %esi
	subq	$1072, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1064(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rax, 392(%rsp)
	call	cudaGetDeviceProperties_v2@PLT
	movl	396(%rsp), %r8d
	xorl	%eax, %eax
	movl	420(%rsp), %edx
	movl	392(%rsp), %ecx
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$44, %edi
	call	malloc@PLT
	movl	$11, %ecx
	movl	$44, %esi
	movss	.LC24(%rip), %xmm0
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	movss	%xmm0, 40(%rbx)
	movq	%rsp, %rdi
	movss	%xmm0, 32(%rbx)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$44, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$67108865, %edi
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$6, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L273
	movq	(%rsp), %rdi
	call	_ZL35__device_stub__Z10testKernelIfEvPT_Pf
.L273:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$44, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$16384, %esi
	movq	%rbx, %rdi
	call	_Z11computeGoldIfEbPT_i
	movq	%rbx, %rdi
	movb	%al, testResult(%rip)
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	1064(%rsp), %rax
	subq	%fs:40, %rax
	je	.L274
	call	__stack_chk_fail@PLT
.L274:
	addq	$1072, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2404:
	.size	_Z7runTestIfEvv, .-_Z7runTestIfEvv
	.section	.text._Z11computeGoldIdEbPT_i,"axG",@progbits,_Z11computeGoldIdEbPT_i,comdat
	.weak	_Z11computeGoldIdEbPT_i
	.type	_Z11computeGoldIdEbPT_i, @function
_Z11computeGoldIdEbPT_i:
.LFB2456:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%eax, %eax
	movl	%esi, %ebp
	xorps	%xmm0, %xmm0
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movsd	.LC29(%rip), %xmm1
.L277:
	cmpl	%ebp, %eax
	jge	.L333
	addsd	%xmm1, %xmm0
	incl	%eax
	jmp	.L277
.L333:
	ucomisd	(%rbx), %xmm0
	jp	.L318
	je	.L308
.L318:
	leaq	.LC5(%rip), %rsi
	jmp	.L332
.L308:
	movsd	.LC29(%rip), %xmm1
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L279:
	cmpl	%ebp, %eax
	jge	.L334
	subsd	%xmm1, %xmm0
	incl	%eax
	jmp	.L279
.L334:
	ucomisd	8(%rbx), %xmm0
	jp	.L319
	xorps	%xmm0, %xmm0
	jne	.L319
	cvtsi2sdl	%ebp, %xmm2
	movsd	.LC30(%rip), %xmm1
	jmp	.L283
.L319:
	leaq	.LC6(%rip), %rsi
.L332:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%eax, %eax
	jmp	.L276
.L288:
	ucomisd	16(%rbx), %xmm0
	jnp	.L335
.L320:
	addsd	%xmm1, %xmm0
.L283:
	comisd	%xmm0, %xmm2
	ja	.L288
	jmp	.L336
.L335:
	jne	.L320
	movsd	.LC26(%rip), %xmm0
	xorps	%xmm1, %xmm1
.L286:
	comisd	%xmm1, %xmm2
	jbe	.L337
	movsd	%xmm2, 8(%rsp)
	movsd	%xmm1, (%rsp)
	call	fmax@PLT
	movsd	(%rsp), %xmm1
	movsd	8(%rsp), %xmm2
	addsd	.LC30(%rip), %xmm1
	jmp	.L286
.L337:
	ucomisd	24(%rbx), %xmm0
	jp	.L321
	je	.L311
.L321:
	leaq	.LC7(%rip), %rsi
	jmp	.L332
.L311:
	movsd	.LC27(%rip), %xmm0
	xorps	%xmm1, %xmm1
.L291:
	comisd	%xmm1, %xmm2
	jbe	.L338
	movsd	%xmm2, 8(%rsp)
	movsd	%xmm1, (%rsp)
	call	fmin@PLT
	movsd	(%rsp), %xmm1
	movsd	8(%rsp), %xmm2
	addsd	.LC30(%rip), %xmm1
	jmp	.L291
.L338:
	ucomisd	32(%rbx), %xmm0
	jp	.L322
	je	.L312
.L322:
	leaq	.LC8(%rip), %rsi
	jmp	.L332
.L312:
	movsd	.LC31(%rip), %xmm1
	movsd	.LC30(%rip), %xmm3
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L294:
	cmpl	%ebp, %eax
	jge	.L339
	comisd	%xmm1, %xmm0
	jnb	.L313
	addsd	%xmm3, %xmm0
	jmp	.L296
.L313:
	xorps	%xmm0, %xmm0
.L296:
	incl	%eax
	jmp	.L294
.L339:
	ucomisd	40(%rbx), %xmm0
	jp	.L323
	je	.L314
.L323:
	leaq	.LC9(%rip), %rsi
	jmp	.L332
.L314:
	xorps	%xmm0, %xmm0
	movsd	.LC30(%rip), %xmm3
	xorl	%eax, %eax
	movaps	%xmm0, %xmm1
.L298:
	cmpl	%ebp, %eax
	jge	.L340
	ucomisd	%xmm1, %xmm0
	jp	.L324
	je	.L315
.L324:
	subsd	%xmm3, %xmm0
	jmp	.L300
.L315:
	movsd	.LC28(%rip), %xmm0
.L300:
	incl	%eax
	jmp	.L298
.L340:
	ucomisd	48(%rbx), %xmm0
	jp	.L325
	movsd	.LC30(%rip), %xmm1
	xorps	%xmm0, %xmm0
	je	.L303
.L325:
	leaq	.LC10(%rip), %rsi
	jmp	.L332
.L306:
	ucomisd	56(%rbx), %xmm0
	jnp	.L341
.L326:
	addsd	%xmm1, %xmm0
.L303:
	comisd	%xmm0, %xmm2
	ja	.L306
	jmp	.L342
.L341:
	jne	.L326
	movb	$1, %al
	jmp	.L276
.L342:
	leaq	.LC14(%rip), %rsi
	jmp	.L332
.L336:
	leaq	.LC15(%rip), %rsi
	jmp	.L332
.L276:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2456:
	.size	_Z11computeGoldIdEbPT_i, .-_Z11computeGoldIdEbPT_i
	.section	.text._Z7runTestIdEvv,"axG",@progbits,_Z7runTestIdEvv,comdat
	.weak	_Z7runTestIdEvv
	.type	_Z7runTestIdEvv, @function
_Z7runTestIdEvv:
.LFB2405:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%esi, %esi
	subq	$1072, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1064(%rsp)
	xorl	%eax, %eax
	leaq	32(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rax, 392(%rsp)
	call	cudaGetDeviceProperties_v2@PLT
	movl	396(%rsp), %r8d
	xorl	%eax, %eax
	movl	420(%rsp), %edx
	movl	392(%rsp), %ecx
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$88, %edi
	call	malloc@PLT
	movl	$22, %ecx
	movl	$88, %esi
	movsd	.LC32(%rip), %xmm0
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	movsd	%xmm0, 80(%rbx)
	movq	%rsp, %rdi
	movsd	%xmm0, 64(%rbx)
	call	cudaMalloc@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$88, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$67108865, %edi
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$6, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L344
	movq	(%rsp), %rdi
	call	_ZL35__device_stub__Z10testKernelIdEvPT_Pd
.L344:
	movq	(%rsp), %rsi
	movl	$2, %ecx
	movl	$88, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$16384, %esi
	movq	%rbx, %rdi
	call	_Z11computeGoldIdEbPT_i
	movq	%rbx, %rdi
	movb	%al, testResult(%rip)
	call	free@PLT
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	1064(%rsp), %rax
	subq	%fs:40, %rax
	je	.L345
	call	__stack_chk_fail@PLT
.L345:
	addq	$1072, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2405:
	.size	_Z7runTestIdEvv, .-_Z7runTestIdEvv
	.section	.rodata.str1.1
.LC33:
	.string	"OK"
.LC34:
	.string	"ERROR!"
.LC35:
	.string	"%s starting...\n"
.LC36:
	.string	"%s completed, returned %s\n"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB2319:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	leaq	.LC35(%rip), %rsi
	movl	$2, %edi
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	sampleName(%rip), %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	call	_Z7runTestIiEvv
	call	_Z7runTestIjEvv
	call	_Z7runTestIyEvv
	call	_Z7runTestIfEvv
	call	_Z7runTestIdEvv
	call	cudaDeviceReset@PLT
	cmpb	$0, testResult(%rip)
	leaq	.LC33(%rip), %rax
	leaq	.LC34(%rip), %rcx
	cmovne	%rax, %rcx
	movq	sampleName(%rip), %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC36(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%edi, %edi
	cmpb	$0, testResult(%rip)
	sete	%dil
	call	exit@PLT
	.cfi_endproc
.LFE2319:
	.size	main, .-main
	.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
	.globl	testResult
	.data
	.type	testResult, @object
	.size	testResult, 1
testResult:
	.byte	1
	.globl	sampleName
	.section	.rodata.str1.1
.LC37:
	.string	"hipSimpleAtomicsTest"
	.section	.data.rel.local,"aw"
	.align 8
	.type	sampleName, @object
	.size	sampleName, 8
sampleName:
	.quad	.LC37
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC18:
	.long	-1015021568
	.align 4
.LC19:
	.long	1132462080
	.align 4
.LC20:
	.long	1124663296
	.align 4
.LC21:
	.long	1092616192
	.align 4
.LC22:
	.long	1065353216
	.align 4
.LC23:
	.long	1099431936
	.align 4
.LC24:
	.long	1132396544
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC26:
	.long	0
	.long	-1066401792
	.align 8
.LC27:
	.long	0
	.long	1081081856
	.align 8
.LC28:
	.long	0
	.long	1080107008
	.align 8
.LC29:
	.long	0
	.long	1076101120
	.align 8
.LC30:
	.long	0
	.long	1072693248
	.align 8
.LC31:
	.long	0
	.long	1076953088
	.align 8
.LC32:
	.long	0
	.long	1081073664
	.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:

// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args "-Xclang" "-fcuda-allow-variadic-functions"

/*
Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
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.
*/

// CHECK: #include "hip/hip_runtime.h"
#include "hip/hip_runtime.h"

#include <algorithm>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <type_traits>

#define EXIT_WAIVED 2

const char* sampleName = "hipSimpleAtomicsTest";

using namespace std;
// Auto-Verification Code
bool testResult = true;

bool computeGoldBitwise(...) {
    return true;
}

template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
bool computeGoldBitwise(T* gpuData, int len) {
    T val = 0xff;
    for (int i = 0; i < len; ++i) {
        // 9th element should be 1
        val &= (2 * i + 7);
    }
    if (val != gpuData[8]) {
        printf("atomicAnd failed\n");
        return false;
    }

    val = 0;
    for (int i = 0; i < len; ++i) {
        // 10th element should be 0xff
        val |= (1 << i);
    }
    if (val != gpuData[9]) {
        printf("atomicOr failed\n");
        return false;
    }

    val = 0xff;
    for (int i = 0; i < len; ++i) {
        // 11th element should be 0xff
        val ^= i;
    }
    if (val != gpuData[10]) {
        printf("atomicXor failed\n");
        return false;
    }

    return true;
}

template<typename T>
bool computeGold(T* gpuData, int len) {
    T val = 0;
    for (int i = 0; i < len; ++i) {
        val += 10;
    }
    if (val != gpuData[0]) {
        printf("atomicAdd failed\n");
        return false;
    }

    val = 0;
    for (int i = 0; i < len; ++i) {
        val -= 10;
    }
    if (val != gpuData[1]) {
        printf("atomicSub failed\n");
        return false;
    }

    bool found = false;
    for (T i = 0; i < len; ++i) {
        // third element should be a member of [0, len)
        if (i == gpuData[2]) {
            found = true;
            break;
        }
    }
    if (!found) {
        printf("atomicExch failed\n");
        return false;
    }

    val = -(1 << 8);
    for (T i = 0; i < len; ++i) {
        // fourth element should be len-1
        val = max(val, i);
    }
    if (val != gpuData[3]) {
        printf("atomicMax failed\n");
        return false;
    }

    val = 1 << 8;
    for (T i = 0; i < len; ++i) {
        val = min(val, i);
    }
    if (val != gpuData[4]) {
        printf("atomicMin failed\n");
        return false;
    }

    int limit = 17;
    val = 0;
    for (int i = 0; i < len; ++i) {
        val = (val >= limit) ? 0 : val + 1;
    }
    if (val != gpuData[5]) {
        printf("atomicInc failed\n");
        return false;
    }

    limit = 137;
    val = 0;
    for (int i = 0; i < len; ++i) {
        val = ((val == 0) || (val > limit)) ? limit : val - 1;
    }
    if (val != gpuData[6]) {
        printf("atomicDec failed\n");
        return false;
    }

    found = false;
    for (T i = 0; i < len; ++i) {
        // eighth element should be a member of [0, len)
        if (i == gpuData[7]) {
            found = true;
            break;
        }
    }
    if (!found) {
        printf("atomicCAS failed\n");
        return false;
    }

    return computeGoldBitwise(gpuData, len);
}

__device__
void testKernelExch(...) {}

template<typename T, typename enable_if<!is_same<T, double>{}>::type* = nullptr>
__device__
void testKernelExch(T* g_odata) {
    // access thread id
    const T tid = blockDim.x * blockIdx.x + threadIdx.x;
    // Atomic exchange
    atomicExch(&g_odata[2], tid);
}

__device__
void testKernelSub(...) {}

template<
    typename T,
    typename enable_if<
        is_same<T, int>{} || is_same<T, unsigned int>{}>::type* = nullptr>
__device__
void testKernelSub(T* g_odata) {
    // Atomic subtraction (final should be 0)
    atomicSub(&g_odata[1], 10);
}

__device__
void testKernelIntegral(...) {}

template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
__device__
void testKernelIntegral(T* g_odata) {
    // access thread id
    const T tid = blockDim.x * blockIdx.x + threadIdx.x;
    // Atomic maximum
    atomicMax(&g_odata[3], tid);
    // Atomic minimum
    atomicMin(&g_odata[4], tid);
    // Atomic increment (modulo 17+1)
    atomicInc((unsigned int*)&g_odata[5], 17);
    // Atomic decrement
    atomicDec((unsigned int*)&g_odata[6], 137);
    // Atomic compare-and-swap
    atomicCAS(&g_odata[7], tid - 1, tid);
    // Atomic AND
    atomicAnd(&g_odata[8], 2 * tid + 7);
    // Atomic OR
    atomicOr(&g_odata[9], 1 << tid);
    // Atomic XOR
    atomicXor(&g_odata[10], tid);
    testKernelSub(g_odata);
}

template<typename T>
__global__ void testKernel(T* g_odata) {
    // Atomic addition
    atomicAdd(&g_odata[0], 10);
    testKernelIntegral(g_odata);
    testKernelExch(g_odata);
}

template<typename T>
void runTest() {
    // CHECK: hipDeviceProp_t deviceProp;
    hipDeviceProp_t deviceProp;
    deviceProp.major = 0;
    deviceProp.minor = 0;
    int dev = 0;
    // CHECK: hipGetDeviceProperties(&deviceProp, dev);
    hipGetDeviceProperties(&deviceProp, dev);
    // Statistics about the GPU device
    printf(
        "> GPU device has %d Multi-Processors, "
        "SM %d.%d compute capabilities\n\n",
        deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
    unsigned int numThreads = 256;
    unsigned int numBlocks = 64;
    unsigned int numData = 11;
    unsigned int memSize = sizeof(T) * numData;

    // Allocate mem for the result on host side
    T* hOData = (T*)malloc(memSize);
    // Initialize the memory
    for (unsigned int i = 0; i < numData; i++) hOData[i] = 0;
    // To make the AND and XOR tests generate something other than 0...
    hOData[8] = hOData[10] = 0xff;
    // Allocate device memory for result
    T* dOData;
    // CHECK: hipMalloc((void**)&dOData, memSize);
    hipMalloc((void**)&dOData, memSize);
    // Copy host memory to device to initialize to zero
    // CHECK: hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice);
    hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice);
    // Execute the kernel
    // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(testKernel), dim3(numBlocks), dim3(numThreads), 0, 0, dOData);
    testKernel<<<numBlocks, numThreads>>>(dOData);
    // Copy result from device to host
    // CHECK: hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost);
    hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost);
    // Compute reference solution
    testResult = computeGold(hOData, numThreads * numBlocks);
    // Cleanup memory
    free(hOData);
    // CHECK: hipFree(dOData);
    hipFree(dOData);
}

int main(int argc, char** argv) {
    printf("%s starting...\n", sampleName);
    runTest<int>();
    runTest<unsigned int>();
    runTest<unsigned long long>();
    runTest<float>();
    runTest<double>();
    // CHECK: hipDeviceReset();
    hipDeviceReset();
    printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!");
    exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
}

	.text
	.file	"atomics.hip"
	.globl	_Z18computeGoldBitwisez         # -- Begin function _Z18computeGoldBitwisez
	.type	_Z18computeGoldBitwisez,@function
_Z18computeGoldBitwisez:                # @_Z18computeGoldBitwisez
	.cfi_startproc
# %bb.0:
	movb	$1, %al
	retq
.Lfunc_end0:
	.size	_Z18computeGoldBitwisez, .Lfunc_end0-_Z18computeGoldBitwisez
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	sampleName(%rip), %rsi
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
	callq	_Z7runTestIiEvv
	callq	_Z7runTestIjEvv
	callq	_Z7runTestIyEvv
	callq	_Z7runTestIfEvv
	callq	_Z7runTestIdEvv
	callq	hipDeviceReset
	movq	sampleName(%rip), %rsi
	cmpb	$0, testResult(%rip)
	movl	$.L.str.4, %eax
	movl	$.L.str.3, %edx
	cmoveq	%rax, %rdx
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	movzbl	testResult(%rip), %edi
	xorl	$1, %edi
	callq	exit
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.section	.text._Z7runTestIiEvv,"axG",@progbits,_Z7runTestIiEvv,comdat
	.weak	_Z7runTestIiEvv                 # -- Begin function _Z7runTestIiEvv
	.type	_Z7runTestIiEvv,@function
_Z7runTestIiEvv:                        # @_Z7runTestIiEvv
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 1504
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	8(%rsp), %rbx
	movq	$0, 360(%rbx)
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	388(%rbx), %esi
	movl	360(%rbx), %edx
	movl	364(%rbx), %ecx
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	movl	$44, %esi
	callq	calloc@PLT
	movq	%rax, %rbx
	movl	$255, %eax
	movl	%eax, 40(%rbx)
	movl	%eax, 32(%rbx)
	movq	%rsp, %r14
	movl	$44, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$44, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967360, %rdi               # imm = 0x100000040
	leaq	192(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_2
# %bb.1:
	movq	(%rsp), %rdi
	callq	_Z25__device_stub__testKernelIiEvPT_
.LBB2_2:
	movq	(%rsp), %rsi
	movl	$44, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	movl	$16384, %esi                    # imm = 0x4000
	callq	_Z11computeGoldIiEbPT_i
	movb	%al, testResult(%rip)
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z7runTestIiEvv, .Lfunc_end2-_Z7runTestIiEvv
	.cfi_endproc
                                        # -- End function
	.section	.text._Z7runTestIjEvv,"axG",@progbits,_Z7runTestIjEvv,comdat
	.weak	_Z7runTestIjEvv                 # -- Begin function _Z7runTestIjEvv
	.type	_Z7runTestIjEvv,@function
_Z7runTestIjEvv:                        # @_Z7runTestIjEvv
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 1504
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	8(%rsp), %rbx
	movq	$0, 360(%rbx)
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	388(%rbx), %esi
	movl	360(%rbx), %edx
	movl	364(%rbx), %ecx
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	movl	$44, %esi
	callq	calloc@PLT
	movq	%rax, %rbx
	movl	$255, %eax
	movl	%eax, 40(%rbx)
	movl	%eax, 32(%rbx)
	movq	%rsp, %r14
	movl	$44, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$44, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967360, %rdi               # imm = 0x100000040
	leaq	192(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_2
# %bb.1:
	movq	(%rsp), %rdi
	callq	_Z25__device_stub__testKernelIjEvPT_
.LBB3_2:
	movq	(%rsp), %rsi
	movl	$44, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	movl	$16384, %esi                    # imm = 0x4000
	callq	_Z11computeGoldIjEbPT_i
	movb	%al, testResult(%rip)
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z7runTestIjEvv, .Lfunc_end3-_Z7runTestIjEvv
	.cfi_endproc
                                        # -- End function
	.section	.text._Z7runTestIyEvv,"axG",@progbits,_Z7runTestIyEvv,comdat
	.weak	_Z7runTestIyEvv                 # -- Begin function _Z7runTestIyEvv
	.type	_Z7runTestIyEvv,@function
_Z7runTestIyEvv:                        # @_Z7runTestIyEvv
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 1504
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	8(%rsp), %rbx
	movq	$0, 360(%rbx)
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	388(%rbx), %esi
	movl	360(%rbx), %edx
	movl	364(%rbx), %ecx
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	movl	$88, %esi
	callq	calloc@PLT
	movq	%rax, %rbx
	movl	$255, %eax
	movq	%rax, 80(%rbx)
	movq	%rax, 64(%rbx)
	movq	%rsp, %r14
	movl	$88, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$88, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967360, %rdi               # imm = 0x100000040
	leaq	192(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_2
# %bb.1:
	movq	(%rsp), %rdi
	callq	_Z25__device_stub__testKernelIyEvPT_
.LBB4_2:
	movq	(%rsp), %rsi
	movl	$88, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	movl	$16384, %esi                    # imm = 0x4000
	callq	_Z11computeGoldIyEbPT_i
	movb	%al, testResult(%rip)
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z7runTestIyEvv, .Lfunc_end4-_Z7runTestIyEvv
	.cfi_endproc
                                        # -- End function
	.section	.text._Z7runTestIfEvv,"axG",@progbits,_Z7runTestIfEvv,comdat
	.weak	_Z7runTestIfEvv                 # -- Begin function _Z7runTestIfEvv
	.type	_Z7runTestIfEvv,@function
_Z7runTestIfEvv:                        # @_Z7runTestIfEvv
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 1504
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	8(%rsp), %rbx
	movq	$0, 360(%rbx)
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	388(%rbx), %esi
	movl	360(%rbx), %edx
	movl	364(%rbx), %ecx
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	movl	$44, %esi
	callq	calloc@PLT
	movq	%rax, %rbx
	movl	$1132396544, %eax               # imm = 0x437F0000
	movl	%eax, 40(%rbx)
	movl	%eax, 32(%rbx)
	movq	%rsp, %r14
	movl	$44, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$44, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967360, %rdi               # imm = 0x100000040
	leaq	192(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_2
# %bb.1:
	movq	(%rsp), %rdi
	callq	_Z25__device_stub__testKernelIfEvPT_
.LBB5_2:
	movq	(%rsp), %rsi
	movl	$44, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	movl	$16384, %esi                    # imm = 0x4000
	callq	_Z11computeGoldIfEbPT_i
	movb	%al, testResult(%rip)
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z7runTestIfEvv, .Lfunc_end5-_Z7runTestIfEvv
	.cfi_endproc
                                        # -- End function
	.section	.text._Z7runTestIdEvv,"axG",@progbits,_Z7runTestIdEvv,comdat
	.weak	_Z7runTestIdEvv                 # -- Begin function _Z7runTestIdEvv
	.type	_Z7runTestIdEvv,@function
_Z7runTestIdEvv:                        # @_Z7runTestIdEvv
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 1504
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	8(%rsp), %rbx
	movq	$0, 360(%rbx)
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipGetDevicePropertiesR0600
	movl	388(%rbx), %esi
	movl	360(%rbx), %edx
	movl	364(%rbx), %ecx
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	movl	$88, %esi
	callq	calloc@PLT
	movq	%rax, %rbx
	movabsq	$4643176031446892544, %rax      # imm = 0x406FE00000000000
	movq	%rax, 80(%rbx)
	movq	%rax, 64(%rbx)
	movq	%rsp, %r14
	movl	$88, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$88, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967360, %rdi               # imm = 0x100000040
	leaq	192(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_2
# %bb.1:
	movq	(%rsp), %rdi
	callq	_Z25__device_stub__testKernelIdEvPT_
.LBB6_2:
	movq	(%rsp), %rsi
	movl	$88, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%rbx, %rdi
	movl	$16384, %esi                    # imm = 0x4000
	callq	_Z11computeGoldIdEbPT_i
	movb	%al, testResult(%rip)
	movq	%rbx, %rdi
	callq	free
	movq	(%rsp), %rdi
	callq	hipFree
	addq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z7runTestIdEvv, .Lfunc_end6-_Z7runTestIdEvv
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelIiEvPT_,"axG",@progbits,_Z25__device_stub__testKernelIiEvPT_,comdat
	.weak	_Z25__device_stub__testKernelIiEvPT_ # -- Begin function _Z25__device_stub__testKernelIiEvPT_
	.type	_Z25__device_stub__testKernelIiEvPT_,@function
_Z25__device_stub__testKernelIiEvPT_:   # @_Z25__device_stub__testKernelIiEvPT_
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z10testKernelIiEvPT_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end7:
	.size	_Z25__device_stub__testKernelIiEvPT_, .Lfunc_end7-_Z25__device_stub__testKernelIiEvPT_
	.cfi_endproc
                                        # -- End function
	.section	.text._Z11computeGoldIiEbPT_i,"axG",@progbits,_Z11computeGoldIiEbPT_i,comdat
	.weak	_Z11computeGoldIiEbPT_i         # -- Begin function _Z11computeGoldIiEbPT_i
	.type	_Z11computeGoldIiEbPT_i,@function
_Z11computeGoldIiEbPT_i:                # @_Z11computeGoldIiEbPT_i
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
                                        # kill: def $esi killed $esi def $rsi
	testl	%esi, %esi
	jle	.LBB8_4
# %bb.1:                                # %._crit_edge
	leal	(%rsi,%rsi), %eax
	leal	(%rax,%rax,4), %ecx
	cmpl	(%rdi), %ecx
	jne	.LBB8_5
# %bb.2:                                # %._crit_edge97.thread
	leal	(%rax,%rax,4), %eax
	addl	4(%rdi), %eax
	jne	.LBB8_3
# %bb.6:
	leal	-1(%rsi), %eax
	cmpl	%eax, 8(%rdi)
	ja	.LBB8_21
# %bb.7:                                # %.lr.ph103.preheader
	movl	$-256, %eax
	xorl	%ecx, %ecx
.LBB8_8:                                # %.lr.ph103
                                        # =>This Inner Loop Header: Depth=1
	cmpl	%ecx, %eax
	cmovlel	%ecx, %eax
	incl	%ecx
	cmpl	%ecx, %esi
	jne	.LBB8_8
# %bb.9:                                # %._crit_edge104
	cmpl	12(%rdi), %eax
	jne	.LBB8_27
# %bb.10:                               # %.lr.ph108.preheader
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	movl	$256, %ecx                      # imm = 0x100
	xorl	%edx, %edx
.LBB8_11:                               # %.lr.ph108
                                        # =>This Inner Loop Header: Depth=1
	cmpl	%edx, %ecx
	cmovgel	%edx, %ecx
	incl	%edx
	cmpl	%edx, %eax
	jne	.LBB8_11
# %bb.12:                               # %._crit_edge109
	cmpl	16(%rdi), %ecx
	jne	.LBB8_22
# %bb.13:                               # %.lr.ph113.preheader
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB8_14:                               # %.lr.ph113
                                        # =>This Inner Loop Header: Depth=1
	cmpl	$17, %edx
	leal	1(%rdx), %edx
                                        # kill: def $edx killed $edx def $rdx
	cmovgel	%ecx, %edx
	decl	%eax
	jne	.LBB8_14
# %bb.15:                               # %._crit_edge114
	cmpl	20(%rdi), %edx
	jne	.LBB8_23
# %bb.16:                               # %.lr.ph118.preheader
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	xorl	%ecx, %ecx
	movl	$137, %edx
.LBB8_17:                               # %.lr.ph118
                                        # =>This Inner Loop Header: Depth=1
	subl	$1, %ecx
	cmovbl	%edx, %ecx
	decl	%eax
	jne	.LBB8_17
# %bb.18:                               # %._crit_edge119
	cmpl	24(%rdi), %ecx
	jne	.LBB8_19
# %bb.24:
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	decl	%eax
	cmpl	%eax, 28(%rdi)
	jbe	.LBB8_26
# %bb.25:                               # %.critedge85
	movl	$.Lstr.42, %edi
	jmp	.LBB8_28
.LBB8_4:                                # %._crit_edge.thread
	cmpl	$0, (%rdi)
	je	.LBB8_20
.LBB8_5:
	movl	$.Lstr.48, %edi
	jmp	.LBB8_28
.LBB8_20:                               # %._crit_edge97
	cmpl	$0, 4(%rdi)
	je	.LBB8_21
.LBB8_3:
	movl	$.Lstr.47, %edi
	jmp	.LBB8_28
.LBB8_21:                               # %.critedge
	movl	$.Lstr.41, %edi
	jmp	.LBB8_28
.LBB8_27:
	movl	$.Lstr.46, %edi
	jmp	.LBB8_28
.LBB8_22:
	movl	$.Lstr.45, %edi
	jmp	.LBB8_28
.LBB8_23:
	movl	$.Lstr.44, %edi
	jmp	.LBB8_28
.LBB8_19:
	movl	$.Lstr.43, %edi
.LBB8_28:
	callq	puts@PLT
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.LBB8_26:
	.cfi_def_cfa_offset 16
                                        # kill: def $esi killed $esi killed $rsi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i # TAILCALL
.Lfunc_end8:
	.size	_Z11computeGoldIiEbPT_i, .Lfunc_end8-_Z11computeGoldIiEbPT_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,"axG",@progbits,_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,comdat
	.weak	_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i # -- Begin function _Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.type	_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,@function
_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i: # @_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB9_13
# %bb.1:                                # %.lr.ph.preheader
	movl	$255, %eax
	movl	$7, %ecx
	movl	%esi, %edx
.LBB9_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	andl	%ecx, %eax
	addl	$2, %ecx
	decl	%edx
	jne	.LBB9_2
# %bb.3:                                # %._crit_edge
	movl	$.Lstr.32, %ecx
	cmpl	32(%rdi), %eax
	jne	.LBB9_14
# %bb.4:                                # %.preheader30
	testl	%esi, %esi
	jle	.LBB9_5
# %bb.11:                               # %.lr.ph35.preheader
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.LBB9_12:                               # %.lr.ph35
                                        # =>This Inner Loop Header: Depth=1
	btsl	%ecx, %eax
	incl	%ecx
	cmpl	%ecx, %esi
	jne	.LBB9_12
	jmp	.LBB9_6
.LBB9_13:                               # %._crit_edge.thread
	movl	$.Lstr.32, %ecx
	xorl	%eax, %eax
	cmpl	$255, 32(%rdi)
	je	.LBB9_6
	jmp	.LBB9_14
.LBB9_5:
	xorl	%eax, %eax
.LBB9_6:                                # %._crit_edge36
	movl	$.Lstr.31, %ecx
	cmpl	36(%rdi), %eax
	jne	.LBB9_14
# %bb.7:                                # %.preheader
	movl	$255, %edx
	testl	%esi, %esi
	jle	.LBB9_10
# %bb.8:                                # %.lr.ph40.preheader
	xorl	%eax, %eax
.LBB9_9:                                # %.lr.ph40
                                        # =>This Inner Loop Header: Depth=1
	xorl	%eax, %edx
	incl	%eax
	cmpl	%eax, %esi
	jne	.LBB9_9
.LBB9_10:                               # %._crit_edge41
	movl	$.Lstr.30, %ecx
	movb	$1, %al
	cmpl	40(%rdi), %edx
	je	.LBB9_15
.LBB9_14:                               # %.sink.split
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rcx, %rdi
	callq	puts@PLT
	xorl	%eax, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB9_15:
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end9:
	.size	_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i, .Lfunc_end9-_Z18computeGoldBitwiseIiTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelIjEvPT_,"axG",@progbits,_Z25__device_stub__testKernelIjEvPT_,comdat
	.weak	_Z25__device_stub__testKernelIjEvPT_ # -- Begin function _Z25__device_stub__testKernelIjEvPT_
	.type	_Z25__device_stub__testKernelIjEvPT_,@function
_Z25__device_stub__testKernelIjEvPT_:   # @_Z25__device_stub__testKernelIjEvPT_
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z10testKernelIjEvPT_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end10:
	.size	_Z25__device_stub__testKernelIjEvPT_, .Lfunc_end10-_Z25__device_stub__testKernelIjEvPT_
	.cfi_endproc
                                        # -- End function
	.section	.text._Z11computeGoldIjEbPT_i,"axG",@progbits,_Z11computeGoldIjEbPT_i,comdat
	.weak	_Z11computeGoldIjEbPT_i         # -- Begin function _Z11computeGoldIjEbPT_i
	.type	_Z11computeGoldIjEbPT_i,@function
_Z11computeGoldIjEbPT_i:                # @_Z11computeGoldIjEbPT_i
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
                                        # kill: def $esi killed $esi def $rsi
	testl	%esi, %esi
	jle	.LBB11_3
# %bb.1:                                # %._crit_edge
	leal	(%rsi,%rsi), %eax
	leal	(%rax,%rax,4), %ecx
	cmpl	(%rdi), %ecx
	jne	.LBB11_30
# %bb.2:                                # %.lr.ph97.preheader
	leal	(%rax,%rax,4), %eax
	negl	%eax
	jmp	.LBB11_5
.LBB11_3:                               # %._crit_edge.thread
	cmpl	$0, (%rdi)
	je	.LBB11_4
.LBB11_30:
	movl	$.Lstr.48, %edi
	jmp	.LBB11_31
.LBB11_4:
	xorl	%eax, %eax
.LBB11_5:                               # %._crit_edge98
	cmpl	4(%rdi), %eax
	jne	.LBB11_32
# %bb.6:                                # %.preheader91
	testl	%esi, %esi
	je	.LBB11_33
# %bb.7:
	leal	-1(%rsi), %eax
	cmpl	%eax, 8(%rdi)
	jbe	.LBB11_8
.LBB11_33:                              # %.critedge
	movl	$.Lstr.41, %edi
	jmp	.LBB11_31
.LBB11_32:
	movl	$.Lstr.47, %edi
.LBB11_31:
	callq	puts@PLT
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.LBB11_8:                               # %.preheader90.preheader
	.cfi_def_cfa_offset 16
	movl	$-256, %ecx
	xorl	%edx, %edx
.LBB11_9:                               # %.preheader90
                                        # =>This Inner Loop Header: Depth=1
	cmpl	%edx, %ecx
	cmovbel	%edx, %ecx
	incl	%edx
	cmpl	%edx, %esi
	jne	.LBB11_9
# %bb.10:
	cmpl	12(%rdi), %ecx
	jne	.LBB11_34
# %bb.11:                               # %.preheader89.preheader
	movl	$256, %ecx                      # imm = 0x100
	xorl	%edx, %edx
.LBB11_12:                              # %.preheader89
                                        # =>This Inner Loop Header: Depth=1
	cmpl	%edx, %ecx
	cmovael	%edx, %ecx
	incl	%edx
	cmpl	%edx, %esi
	jne	.LBB11_12
# %bb.13:
	cmpl	16(%rdi), %ecx
	jne	.LBB11_19
# %bb.14:                               # %.preheader88
	testl	%esi, %esi
	jle	.LBB11_20
# %bb.15:                               # %.lr.ph108.preheader
	xorl	%ecx, %ecx
	movl	%esi, %edx
	xorl	%r8d, %r8d
.LBB11_16:                              # %.lr.ph108
                                        # =>This Inner Loop Header: Depth=1
	cmpl	$17, %r8d
	leal	1(%r8), %r8d
                                        # kill: def $r8d killed $r8d def $r8
	cmovael	%ecx, %r8d
	decl	%edx
	jne	.LBB11_16
# %bb.17:                               # %._crit_edge109
	cmpl	20(%rdi), %r8d
	jne	.LBB11_18
# %bb.24:                               # %.preheader87
	testl	%esi, %esi
	jle	.LBB11_21
# %bb.25:                               # %.lr.ph113.preheader
	xorl	%ecx, %ecx
	movl	$137, %edx
	movl	%esi, %r8d
.LBB11_26:                              # %.lr.ph113
                                        # =>This Inner Loop Header: Depth=1
	leal	-138(%rcx), %r9d
	decl	%ecx
	cmpl	$-137, %r9d
	cmovbl	%edx, %ecx
	decl	%r8d
	jne	.LBB11_26
	jmp	.LBB11_22
.LBB11_34:
	movl	$.Lstr.46, %edi
	jmp	.LBB11_31
.LBB11_19:
	movl	$.Lstr.45, %edi
	jmp	.LBB11_31
.LBB11_20:                              # %._crit_edge109.thread
	cmpl	$0, 20(%rdi)
	je	.LBB11_21
.LBB11_18:
	movl	$.Lstr.44, %edi
	jmp	.LBB11_31
.LBB11_21:
	xorl	%ecx, %ecx
.LBB11_22:                              # %._crit_edge114
	cmpl	24(%rdi), %ecx
	jne	.LBB11_23
# %bb.27:
	cmpl	%eax, 28(%rdi)
	jbe	.LBB11_29
# %bb.28:                               # %.critedge61
	movl	$.Lstr.42, %edi
	jmp	.LBB11_31
.LBB11_23:
	movl	$.Lstr.43, %edi
	jmp	.LBB11_31
.LBB11_29:
                                        # kill: def $esi killed $esi killed $rsi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i # TAILCALL
.Lfunc_end11:
	.size	_Z11computeGoldIjEbPT_i, .Lfunc_end11-_Z11computeGoldIjEbPT_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,"axG",@progbits,_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,comdat
	.weak	_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i # -- Begin function _Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.type	_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,@function
_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i: # @_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB12_13
# %bb.1:                                # %.lr.ph.preheader
	movl	$255, %eax
	movl	$7, %ecx
	movl	%esi, %edx
.LBB12_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	andl	%ecx, %eax
	addl	$2, %ecx
	decl	%edx
	jne	.LBB12_2
# %bb.3:                                # %._crit_edge
	movl	$.Lstr.32, %ecx
	cmpl	32(%rdi), %eax
	jne	.LBB12_14
# %bb.4:                                # %.preheader30
	testl	%esi, %esi
	jle	.LBB12_5
# %bb.11:                               # %.lr.ph35.preheader
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.LBB12_12:                              # %.lr.ph35
                                        # =>This Inner Loop Header: Depth=1
	btsl	%ecx, %eax
	incl	%ecx
	cmpl	%ecx, %esi
	jne	.LBB12_12
	jmp	.LBB12_6
.LBB12_13:                              # %._crit_edge.thread
	movl	$.Lstr.32, %ecx
	xorl	%eax, %eax
	cmpl	$255, 32(%rdi)
	je	.LBB12_6
	jmp	.LBB12_14
.LBB12_5:
	xorl	%eax, %eax
.LBB12_6:                               # %._crit_edge36
	movl	$.Lstr.31, %ecx
	cmpl	36(%rdi), %eax
	jne	.LBB12_14
# %bb.7:                                # %.preheader
	movl	$255, %edx
	testl	%esi, %esi
	jle	.LBB12_10
# %bb.8:                                # %.lr.ph40.preheader
	xorl	%eax, %eax
.LBB12_9:                               # %.lr.ph40
                                        # =>This Inner Loop Header: Depth=1
	xorl	%eax, %edx
	incl	%eax
	cmpl	%eax, %esi
	jne	.LBB12_9
.LBB12_10:                              # %._crit_edge41
	movl	$.Lstr.30, %ecx
	movb	$1, %al
	cmpl	40(%rdi), %edx
	je	.LBB12_15
.LBB12_14:                              # %.sink.split
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rcx, %rdi
	callq	puts@PLT
	xorl	%eax, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB12_15:
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end12:
	.size	_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i, .Lfunc_end12-_Z18computeGoldBitwiseIjTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelIyEvPT_,"axG",@progbits,_Z25__device_stub__testKernelIyEvPT_,comdat
	.weak	_Z25__device_stub__testKernelIyEvPT_ # -- Begin function _Z25__device_stub__testKernelIyEvPT_
	.type	_Z25__device_stub__testKernelIyEvPT_,@function
_Z25__device_stub__testKernelIyEvPT_:   # @_Z25__device_stub__testKernelIyEvPT_
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z10testKernelIyEvPT_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end13:
	.size	_Z25__device_stub__testKernelIyEvPT_, .Lfunc_end13-_Z25__device_stub__testKernelIyEvPT_
	.cfi_endproc
                                        # -- End function
	.section	.text._Z11computeGoldIyEbPT_i,"axG",@progbits,_Z11computeGoldIyEbPT_i,comdat
	.weak	_Z11computeGoldIyEbPT_i         # -- Begin function _Z11computeGoldIyEbPT_i
	.type	_Z11computeGoldIyEbPT_i,@function
_Z11computeGoldIyEbPT_i:                # @_Z11computeGoldIyEbPT_i
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
                                        # kill: def $esi killed $esi def $rsi
	testl	%esi, %esi
	jle	.LBB14_3
# %bb.1:                                # %._crit_edge
	leal	-1(%rsi), %eax
	leaq	(%rax,%rax,4), %rcx
	leaq	10(,%rcx,2), %rcx
	cmpq	(%rdi), %rcx
	jne	.LBB14_30
# %bb.2:                                # %.lr.ph99.preheader
	addq	%rax, %rax
	leaq	(%rax,%rax,4), %rcx
	movq	$-10, %rax
	subq	%rcx, %rax
	jmp	.LBB14_5
.LBB14_3:                               # %._crit_edge.thread
	cmpq	$0, (%rdi)
	je	.LBB14_4
.LBB14_30:
	movl	$.Lstr.48, %edi
	jmp	.LBB14_31
.LBB14_4:
	xorl	%eax, %eax
.LBB14_5:                               # %._crit_edge100
	cmpq	8(%rdi), %rax
	jne	.LBB14_32
# %bb.6:                                # %.preheader92
	testl	%esi, %esi
	je	.LBB14_33
# %bb.7:
	movslq	%esi, %rax
	leaq	-1(%rax), %rcx
	cmpq	%rcx, 16(%rdi)
	jbe	.LBB14_8
.LBB14_33:                              # %.critedge
	movl	$.Lstr.41, %edi
	jmp	.LBB14_31
.LBB14_32:
	movl	$.Lstr.47, %edi
.LBB14_31:
	callq	puts@PLT
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.LBB14_8:                               # %.preheader91.preheader
	.cfi_def_cfa_offset 16
	movq	$-256, %rcx
	xorl	%edx, %edx
.LBB14_9:                               # %.preheader91
                                        # =>This Inner Loop Header: Depth=1
	cmpq	%rdx, %rcx
	cmovbeq	%rdx, %rcx
	incq	%rdx
	cmpq	%rdx, %rax
	jne	.LBB14_9
# %bb.10:
	cmpq	24(%rdi), %rcx
	jne	.LBB14_34
# %bb.11:                               # %.preheader90.preheader
	movl	$256, %ecx                      # imm = 0x100
	xorl	%edx, %edx
.LBB14_12:                              # %.preheader90
                                        # =>This Inner Loop Header: Depth=1
	cmpq	%rdx, %rcx
	cmovaeq	%rdx, %rcx
	incq	%rdx
	cmpq	%rdx, %rax
	jne	.LBB14_12
# %bb.13:
	cmpq	32(%rdi), %rcx
	jne	.LBB14_19
# %bb.14:                               # %.preheader89
	testl	%esi, %esi
	jle	.LBB14_20
# %bb.15:                               # %.lr.ph110.preheader
	xorl	%ecx, %ecx
	movl	%esi, %edx
.LBB14_16:                              # %.lr.ph110
                                        # =>This Inner Loop Header: Depth=1
	leaq	1(%rcx), %r8
	cmpq	$17, %rcx
	movl	$0, %ecx
	cmovbq	%r8, %rcx
	decl	%edx
	jne	.LBB14_16
# %bb.17:                               # %._crit_edge111
	cmpq	40(%rdi), %rcx
	jne	.LBB14_18
# %bb.24:                               # %.preheader88
	testl	%esi, %esi
	jle	.LBB14_21
# %bb.25:                               # %.lr.ph115.preheader
	xorl	%ecx, %ecx
	movl	$137, %edx
	movl	%esi, %r8d
.LBB14_26:                              # %.lr.ph115
                                        # =>This Inner Loop Header: Depth=1
	leaq	-138(%rcx), %r9
	decq	%rcx
	cmpq	$-137, %r9
	cmovbq	%rdx, %rcx
	decl	%r8d
	jne	.LBB14_26
	jmp	.LBB14_22
.LBB14_34:
	movl	$.Lstr.46, %edi
	jmp	.LBB14_31
.LBB14_19:
	movl	$.Lstr.45, %edi
	jmp	.LBB14_31
.LBB14_20:                              # %._crit_edge111.thread
	cmpq	$0, 40(%rdi)
	je	.LBB14_21
.LBB14_18:
	movl	$.Lstr.44, %edi
	jmp	.LBB14_31
.LBB14_21:
	xorl	%ecx, %ecx
.LBB14_22:                              # %._crit_edge116
	cmpq	48(%rdi), %rcx
	jne	.LBB14_23
# %bb.27:
	cmpq	$1, %rax
	adcq	$-1, %rax
	cmpq	%rax, 56(%rdi)
	jbe	.LBB14_29
# %bb.28:                               # %.critedge62
	movl	$.Lstr.42, %edi
	jmp	.LBB14_31
.LBB14_23:
	movl	$.Lstr.43, %edi
	jmp	.LBB14_31
.LBB14_29:
                                        # kill: def $esi killed $esi killed $rsi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i # TAILCALL
.Lfunc_end14:
	.size	_Z11computeGoldIyEbPT_i, .Lfunc_end14-_Z11computeGoldIyEbPT_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,"axG",@progbits,_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,comdat
	.weak	_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i # -- Begin function _Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.type	_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i,@function
_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i: # @_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB15_7
# %bb.1:                                # %.lr.ph.preheader
	movl	%esi, %eax
	addq	%rax, %rax
	movl	$255, %edx
	xorl	%ecx, %ecx
.LBB15_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	leaq	7(%rcx), %r8
	andq	%r8, %rdx
	addq	$2, %rcx
	cmpq	%rcx, %rax
	jne	.LBB15_2
# %bb.3:                                # %._crit_edge
	movl	$.Lstr.32, %ecx
	cmpq	64(%rdi), %rdx
	jne	.LBB15_15
# %bb.4:                                # %.preheader30
	xorl	%ecx, %ecx
	testl	%esi, %esi
	jle	.LBB15_8
# %bb.5:                                # %.lr.ph35.preheader
	xorl	%eax, %eax
.LBB15_6:                               # %.lr.ph35
                                        # =>This Inner Loop Header: Depth=1
	movl	$1, %edx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%rdx, %rax
	incl	%ecx
	cmpl	%ecx, %esi
	jne	.LBB15_6
	jmp	.LBB15_9
.LBB15_7:                               # %._crit_edge.thread
	movl	$.Lstr.32, %ecx
	cmpq	$255, 64(%rdi)
	jne	.LBB15_15
.LBB15_8:
	xorl	%eax, %eax
.LBB15_9:                               # %._crit_edge36
	movl	$.Lstr.31, %ecx
	cmpq	72(%rdi), %rax
	jne	.LBB15_15
# %bb.10:                               # %.preheader
	testl	%esi, %esi
	jle	.LBB15_11
# %bb.12:                               # %.lr.ph40.preheader
	movl	%esi, %eax
	movl	$255, %edx
	xorl	%ecx, %ecx
.LBB15_13:                              # %.lr.ph40
                                        # =>This Inner Loop Header: Depth=1
	xorq	%rcx, %rdx
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB15_13
	jmp	.LBB15_14
.LBB15_11:
	movl	$255, %edx
.LBB15_14:                              # %._crit_edge41
	movl	$.Lstr.30, %ecx
	movb	$1, %al
	cmpq	80(%rdi), %rdx
	je	.LBB15_16
.LBB15_15:                              # %.sink.split
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rcx, %rdi
	callq	puts@PLT
	xorl	%eax, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB15_16:
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end15:
	.size	_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i, .Lfunc_end15-_Z18computeGoldBitwiseIyTnPNSt9enable_ifIXtlSt11is_integralIT_EEEvE4typeELPv0EEbPS2_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelIfEvPT_,"axG",@progbits,_Z25__device_stub__testKernelIfEvPT_,comdat
	.weak	_Z25__device_stub__testKernelIfEvPT_ # -- Begin function _Z25__device_stub__testKernelIfEvPT_
	.type	_Z25__device_stub__testKernelIfEvPT_,@function
_Z25__device_stub__testKernelIfEvPT_:   # @_Z25__device_stub__testKernelIfEvPT_
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z10testKernelIfEvPT_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end16:
	.size	_Z25__device_stub__testKernelIfEvPT_, .Lfunc_end16-_Z25__device_stub__testKernelIfEvPT_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z11computeGoldIfEbPT_i
.LCPI17_0:
	.long	0x41200000                      # float 10
.LCPI17_1:
	.long	0xc1200000                      # float -10
.LCPI17_2:
	.long	0x3f800000                      # float 1
.LCPI17_3:
	.long	0xc3800000                      # float -256
.LCPI17_4:
	.long	0x43800000                      # float 256
.LCPI17_5:
	.long	0x41880000                      # float 17
.LCPI17_6:
	.long	0xbf800000                      # float -1
.LCPI17_7:
	.long	0x43090000                      # float 137
	.section	.text._Z11computeGoldIfEbPT_i,"axG",@progbits,_Z11computeGoldIfEbPT_i,comdat
	.weak	_Z11computeGoldIfEbPT_i
	.type	_Z11computeGoldIfEbPT_i,@function
_Z11computeGoldIfEbPT_i:                # @_Z11computeGoldIfEbPT_i
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	testl	%esi, %esi
	jle	.LBB17_7
# %bb.1:                                # %.lr.ph.preheader
	xorps	%xmm0, %xmm0
	movss	.LCPI17_0(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	movl	%esi, %eax
.LBB17_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	addss	%xmm1, %xmm0
	decl	%eax
	jne	.LBB17_2
# %bb.3:                                # %._crit_edge
	movl	$.Lstr.48, %eax
	ucomiss	(%rdi), %xmm0
	jne	.LBB17_28
	jp	.LBB17_28
# %bb.4:                                # %.preheader82
	xorps	%xmm0, %xmm0
	testl	%esi, %esi
	jle	.LBB17_8
# %bb.5:                                # %.lr.ph88.preheader
	movss	.LCPI17_1(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
	movl	%esi, %eax
.LBB17_6:                               # %.lr.ph88
                                        # =>This Inner Loop Header: Depth=1
	addss	%xmm1, %xmm0
	decl	%eax
	jne	.LBB17_6
	jmp	.LBB17_8
.LBB17_7:                               # %._crit_edge.thread
	movss	(%rdi), %xmm1                   # xmm1 = mem[0],zero,zero,zero
	xorps	%xmm0, %xmm0
	movl	$.Lstr.48, %eax
	ucomiss	%xmm0, %xmm1
	jne	.LBB17_28
	jp	.LBB17_28
.LBB17_8:                               # %._crit_edge89
	movl	$.Lstr.47, %eax
	ucomiss	4(%rdi), %xmm0
	jne	.LBB17_28
	jnp	.LBB17_9
.LBB17_28:                              # %.loopexit.sink.split
	movq	%rax, %rdi
	callq	puts@PLT
	xorl	%eax, %eax
.LBB17_29:                              # %.loopexit
                                        # kill: def $al killed $al killed $eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.LBB17_9:                               # %.preheader81
	.cfi_def_cfa_offset 16
	movl	$.Lstr.41, %eax
	testl	%esi, %esi
	jle	.LBB17_28
# %bb.10:                               # %.lr.ph92
	xorps	%xmm0, %xmm0
	cvtsi2ss	%esi, %xmm0
	movss	8(%rdi), %xmm2                  # xmm2 = mem[0],zero,zero,zero
	xorps	%xmm3, %xmm3
	movss	.LCPI17_2(%rip), %xmm1          # xmm1 = mem[0],zero,zero,zero
.LBB17_11:                              # =>This Inner Loop Header: Depth=1
	ucomiss	%xmm2, %xmm3
	jne	.LBB17_30
	jnp	.LBB17_12
.LBB17_30:                              #   in Loop: Header=BB17_11 Depth=1
	addss	%xmm1, %xmm3
	ucomiss	%xmm3, %xmm0
	ja	.LBB17_11
	jmp	.LBB17_28
.LBB17_12:                              # %.lr.ph95.preheader
	movss	.LCPI17_3(%rip), %xmm2          # xmm2 = mem[0],zero,zero,zero
	xorps	%xmm3, %xmm3
.LBB17_13:                              # %.lr.ph95
                                        # =>This Inner Loop Header: Depth=1
	movaps	%xmm2, %xmm4
	movaps	%xmm3, %xmm2
	maxss	%xmm4, %xmm2
	addss	%xmm1, %xmm3
	ucomiss	%xmm3, %xmm0
	ja	.LBB17_13
# %bb.14:                               # %._crit_edge96
	movl	$.Lstr.46, %eax
	ucomiss	12(%rdi), %xmm2
	jne	.LBB17_28
	jp	.LBB17_28
# %bb.15:                               # %.lr.ph100.preheader
	movss	.LCPI17_4(%rip), %xmm2          # xmm2 = mem[0],zero,zero,zero
	xorps	%xmm3, %xmm3
.LBB17_16:                              # %.lr.ph100
                                        # =>This Inner Loop Header: Depth=1
	minss	%xmm3, %xmm2
	addss	%xmm1, %xmm3
	ucomiss	%xmm3, %xmm0
	ja	.LBB17_16
# %bb.17:                               # %._crit_edge101
	movl	$.Lstr.45, %eax
	ucomiss	16(%rdi), %xmm2
	jne	.LBB17_28
	jp	.LBB17_28
# %bb.18:                               # %.lr.ph105.preheader
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	xorps	%xmm2, %xmm2
	movss	.LCPI17_5(%rip), %xmm3          # xmm3 = mem[0],zero,zero,zero
.LBB17_19:                              # %.lr.ph105
                                        # =>This Inner Loop Header: Depth=1
	movaps	%xmm3, %xmm4
	cmpless	%xmm2, %xmm4
	addss	%xmm1, %xmm2
	andnps	%xmm2, %xmm4
	movaps	%xmm4, %xmm2
	decl	%eax
	jne	.LBB17_19
# %bb.20:                               # %._crit_edge106
	movl	$.Lstr.44, %eax
	ucomiss	20(%rdi), %xmm2
	jne	.LBB17_28
	jp	.LBB17_28
# %bb.21:                               # %.lr.ph110.preheader
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	xorps	%xmm2, %xmm2
	movss	.LCPI17_6(%rip), %xmm3          # xmm3 = mem[0],zero,zero,zero
	movss	.LCPI17_7(%rip), %xmm5          # xmm5 = mem[0],zero,zero,zero
	xorps	%xmm4, %xmm4
.LBB17_22:                              # %.lr.ph110
                                        # =>This Inner Loop Header: Depth=1
	movaps	%xmm4, %xmm6
	addss	%xmm3, %xmm6
	movaps	%xmm5, %xmm7
	cmpltss	%xmm4, %xmm7
	movaps	%xmm7, %xmm8
	andps	%xmm5, %xmm8
	andnps	%xmm6, %xmm7
	orps	%xmm8, %xmm7
	cmpeqss	%xmm2, %xmm4
	movaps	%xmm4, %xmm6
	andnps	%xmm7, %xmm6
	andps	%xmm5, %xmm4
	orps	%xmm6, %xmm4
	decl	%eax
	jne	.LBB17_22
# %bb.23:                               # %._crit_edge111
	movl	$.Lstr.43, %eax
	ucomiss	24(%rdi), %xmm4
	jne	.LBB17_28
	jp	.LBB17_28
# %bb.24:                               # %.lr.ph114
	movss	28(%rdi), %xmm2                 # xmm2 = mem[0],zero,zero,zero
	xorps	%xmm3, %xmm3
	movl	$.Lstr.42, %eax
.LBB17_25:                              # =>This Inner Loop Header: Depth=1
	ucomiss	%xmm2, %xmm3
	jne	.LBB17_27
	jnp	.LBB17_26
.LBB17_27:                              #   in Loop: Header=BB17_25 Depth=1
	addss	%xmm1, %xmm3
	ucomiss	%xmm3, %xmm0
	ja	.LBB17_25
	jmp	.LBB17_28
.LBB17_26:
	movb	$1, %al
	jmp	.LBB17_29
.Lfunc_end17:
	.size	_Z11computeGoldIfEbPT_i, .Lfunc_end17-_Z11computeGoldIfEbPT_i
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25__device_stub__testKernelIdEvPT_,"axG",@progbits,_Z25__device_stub__testKernelIdEvPT_,comdat
	.weak	_Z25__device_stub__testKernelIdEvPT_ # -- Begin function _Z25__device_stub__testKernelIdEvPT_
	.type	_Z25__device_stub__testKernelIdEvPT_,@function
_Z25__device_stub__testKernelIdEvPT_:   # @_Z25__device_stub__testKernelIdEvPT_
	.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	$64, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsp, %rbx
	movq	%rax, (%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z10testKernelIdEvPT_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	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_end18:
	.size	_Z25__device_stub__testKernelIdEvPT_, .Lfunc_end18-_Z25__device_stub__testKernelIdEvPT_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z11computeGoldIdEbPT_i
.LCPI19_0:
	.quad	0x4024000000000000              # double 10
.LCPI19_1:
	.quad	0xc024000000000000              # double -10
.LCPI19_2:
	.quad	0x3ff0000000000000              # double 1
.LCPI19_3:
	.quad	0xc070000000000000              # double -256
.LCPI19_4:
	.quad	0x4070000000000000              # double 256
.LCPI19_5:
	.quad	0x4031000000000000              # double 17
.LCPI19_6:
	.quad	0xbff0000000000000              # double -1
.LCPI19_7:
	.quad	0x4061200000000000              # double 137
	.section	.text._Z11computeGoldIdEbPT_i,"axG",@progbits,_Z11computeGoldIdEbPT_i,comdat
	.weak	_Z11computeGoldIdEbPT_i
	.type	_Z11computeGoldIdEbPT_i,@function
_Z11computeGoldIdEbPT_i:                # @_Z11computeGoldIdEbPT_i
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	testl	%esi, %esi
	jle	.LBB19_7
# %bb.1:                                # %.lr.ph.preheader
	xorpd	%xmm0, %xmm0
	movsd	.LCPI19_0(%rip), %xmm1          # xmm1 = mem[0],zero
	movl	%esi, %eax
.LBB19_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	addsd	%xmm1, %xmm0
	decl	%eax
	jne	.LBB19_2
# %bb.3:                                # %._crit_edge
	movl	$.Lstr.48, %eax
	ucomisd	(%rdi), %xmm0
	jne	.LBB19_28
	jp	.LBB19_28
# %bb.4:                                # %.preheader82
	xorpd	%xmm0, %xmm0
	testl	%esi, %esi
	jle	.LBB19_8
# %bb.5:                                # %.lr.ph88.preheader
	movsd	.LCPI19_1(%rip), %xmm1          # xmm1 = mem[0],zero
	movl	%esi, %eax
.LBB19_6:                               # %.lr.ph88
                                        # =>This Inner Loop Header: Depth=1
	addsd	%xmm1, %xmm0
	decl	%eax
	jne	.LBB19_6
	jmp	.LBB19_8
.LBB19_7:                               # %._crit_edge.thread
	movsd	(%rdi), %xmm1                   # xmm1 = mem[0],zero
	xorpd	%xmm0, %xmm0
	movl	$.Lstr.48, %eax
	ucomisd	%xmm0, %xmm1
	jne	.LBB19_28
	jp	.LBB19_28
.LBB19_8:                               # %._crit_edge89
	movl	$.Lstr.47, %eax
	ucomisd	8(%rdi), %xmm0
	jne	.LBB19_28
	jnp	.LBB19_9
.LBB19_28:                              # %.loopexit.sink.split
	movq	%rax, %rdi
	callq	puts@PLT
	xorl	%eax, %eax
.LBB19_29:                              # %.loopexit
                                        # kill: def $al killed $al killed $eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.LBB19_9:                               # %.preheader81
	.cfi_def_cfa_offset 16
	movl	$.Lstr.41, %eax
	testl	%esi, %esi
	jle	.LBB19_28
# %bb.10:                               # %.lr.ph92
	xorps	%xmm0, %xmm0
	cvtsi2sd	%esi, %xmm0
	movsd	16(%rdi), %xmm2                 # xmm2 = mem[0],zero
	xorpd	%xmm3, %xmm3
	movsd	.LCPI19_2(%rip), %xmm1          # xmm1 = mem[0],zero
.LBB19_11:                              # =>This Inner Loop Header: Depth=1
	ucomisd	%xmm2, %xmm3
	jne	.LBB19_30
	jnp	.LBB19_12
.LBB19_30:                              #   in Loop: Header=BB19_11 Depth=1
	addsd	%xmm1, %xmm3
	ucomisd	%xmm3, %xmm0
	ja	.LBB19_11
	jmp	.LBB19_28
.LBB19_12:                              # %.lr.ph95.preheader
	movsd	.LCPI19_3(%rip), %xmm2          # xmm2 = mem[0],zero
	xorpd	%xmm3, %xmm3
.LBB19_13:                              # %.lr.ph95
                                        # =>This Inner Loop Header: Depth=1
	movapd	%xmm2, %xmm4
	movapd	%xmm3, %xmm2
	maxsd	%xmm4, %xmm2
	addsd	%xmm1, %xmm3
	ucomisd	%xmm3, %xmm0
	ja	.LBB19_13
# %bb.14:                               # %._crit_edge96
	movl	$.Lstr.46, %eax
	ucomisd	24(%rdi), %xmm2
	jne	.LBB19_28
	jp	.LBB19_28
# %bb.15:                               # %.lr.ph100.preheader
	movsd	.LCPI19_4(%rip), %xmm2          # xmm2 = mem[0],zero
	xorpd	%xmm3, %xmm3
.LBB19_16:                              # %.lr.ph100
                                        # =>This Inner Loop Header: Depth=1
	minsd	%xmm3, %xmm2
	addsd	%xmm1, %xmm3
	ucomisd	%xmm3, %xmm0
	ja	.LBB19_16
# %bb.17:                               # %._crit_edge101
	movl	$.Lstr.45, %eax
	ucomisd	32(%rdi), %xmm2
	jne	.LBB19_28
	jp	.LBB19_28
# %bb.18:                               # %.lr.ph105.preheader
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	xorpd	%xmm2, %xmm2
	movsd	.LCPI19_5(%rip), %xmm3          # xmm3 = mem[0],zero
.LBB19_19:                              # %.lr.ph105
                                        # =>This Inner Loop Header: Depth=1
	movapd	%xmm3, %xmm4
	cmplesd	%xmm2, %xmm4
	addsd	%xmm1, %xmm2
	andnpd	%xmm2, %xmm4
	movapd	%xmm4, %xmm2
	decl	%eax
	jne	.LBB19_19
# %bb.20:                               # %._crit_edge106
	movl	$.Lstr.44, %eax
	ucomisd	40(%rdi), %xmm2
	jne	.LBB19_28
	jp	.LBB19_28
# %bb.21:                               # %.lr.ph110.preheader
	cmpl	$2, %esi
	movl	$1, %eax
	cmovgel	%esi, %eax
	xorpd	%xmm2, %xmm2
	movsd	.LCPI19_6(%rip), %xmm3          # xmm3 = mem[0],zero
	movsd	.LCPI19_7(%rip), %xmm5          # xmm5 = mem[0],zero
	xorpd	%xmm4, %xmm4
.LBB19_22:                              # %.lr.ph110
                                        # =>This Inner Loop Header: Depth=1
	movapd	%xmm4, %xmm6
	addsd	%xmm3, %xmm6
	movapd	%xmm5, %xmm7
	cmpltsd	%xmm4, %xmm7
	movapd	%xmm7, %xmm8
	andpd	%xmm5, %xmm8
	andnpd	%xmm6, %xmm7
	orpd	%xmm8, %xmm7
	cmpeqsd	%xmm2, %xmm4
	movapd	%xmm4, %xmm6
	andnpd	%xmm7, %xmm6
	andpd	%xmm5, %xmm4
	orpd	%xmm6, %xmm4
	decl	%eax
	jne	.LBB19_22
# %bb.23:                               # %._crit_edge111
	movl	$.Lstr.43, %eax
	ucomisd	48(%rdi), %xmm4
	jne	.LBB19_28
	jp	.LBB19_28
# %bb.24:                               # %.lr.ph114
	movsd	56(%rdi), %xmm2                 # xmm2 = mem[0],zero
	xorpd	%xmm3, %xmm3
	movl	$.Lstr.42, %eax
.LBB19_25:                              # =>This Inner Loop Header: Depth=1
	ucomisd	%xmm2, %xmm3
	jne	.LBB19_27
	jnp	.LBB19_26
.LBB19_27:                              #   in Loop: Header=BB19_25 Depth=1
	addsd	%xmm1, %xmm3
	ucomisd	%xmm3, %xmm0
	ja	.LBB19_25
	jmp	.LBB19_28
.LBB19_26:
	movb	$1, %al
	jmp	.LBB19_29
.Lfunc_end19:
	.size	_Z11computeGoldIdEbPT_i, .Lfunc_end19-_Z11computeGoldIdEbPT_i
	.cfi_endproc
                                        # -- End function
	.text
	.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	.LBB20_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB20_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10testKernelIiEvPT_, %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	$_Z10testKernelIjEvPT_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10testKernelIyEvPT_, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10testKernelIfEvPT_, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10testKernelIdEvPT_, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end20:
	.size	__hip_module_ctor, .Lfunc_end20-__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	.LBB21_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
.LBB21_2:
	retq
.Lfunc_end21:
	.size	__hip_module_dtor, .Lfunc_end21-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"hipSimpleAtomicsTest"
	.size	.L.str, 21

	.type	sampleName,@object              # @sampleName
	.data
	.globl	sampleName
	.p2align	3, 0x0
sampleName:
	.quad	.L.str
	.size	sampleName, 8

	.type	testResult,@object              # @testResult
	.globl	testResult
testResult:
	.byte	1                               # 0x1
	.size	testResult, 1

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"%s starting...\n"
	.size	.L.str.1, 16

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"%s completed, returned %s\n"
	.size	.L.str.2, 27

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"OK"
	.size	.L.str.3, 3

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"ERROR!"
	.size	.L.str.4, 7

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"> GPU device has %d Multi-Processors, SM %d.%d compute capabilities\n\n"
	.size	.L.str.5, 70

	.type	_Z10testKernelIiEvPT_,@object   # @_Z10testKernelIiEvPT_
	.section	.rodata._Z10testKernelIiEvPT_,"aG",@progbits,_Z10testKernelIiEvPT_,comdat
	.weak	_Z10testKernelIiEvPT_
	.p2align	3, 0x0
_Z10testKernelIiEvPT_:
	.quad	_Z25__device_stub__testKernelIiEvPT_
	.size	_Z10testKernelIiEvPT_, 8

	.type	_Z10testKernelIjEvPT_,@object   # @_Z10testKernelIjEvPT_
	.section	.rodata._Z10testKernelIjEvPT_,"aG",@progbits,_Z10testKernelIjEvPT_,comdat
	.weak	_Z10testKernelIjEvPT_
	.p2align	3, 0x0
_Z10testKernelIjEvPT_:
	.quad	_Z25__device_stub__testKernelIjEvPT_
	.size	_Z10testKernelIjEvPT_, 8

	.type	_Z10testKernelIyEvPT_,@object   # @_Z10testKernelIyEvPT_
	.section	.rodata._Z10testKernelIyEvPT_,"aG",@progbits,_Z10testKernelIyEvPT_,comdat
	.weak	_Z10testKernelIyEvPT_
	.p2align	3, 0x0
_Z10testKernelIyEvPT_:
	.quad	_Z25__device_stub__testKernelIyEvPT_
	.size	_Z10testKernelIyEvPT_, 8

	.type	_Z10testKernelIfEvPT_,@object   # @_Z10testKernelIfEvPT_
	.section	.rodata._Z10testKernelIfEvPT_,"aG",@progbits,_Z10testKernelIfEvPT_,comdat
	.weak	_Z10testKernelIfEvPT_
	.p2align	3, 0x0
_Z10testKernelIfEvPT_:
	.quad	_Z25__device_stub__testKernelIfEvPT_
	.size	_Z10testKernelIfEvPT_, 8

	.type	_Z10testKernelIdEvPT_,@object   # @_Z10testKernelIdEvPT_
	.section	.rodata._Z10testKernelIdEvPT_,"aG",@progbits,_Z10testKernelIdEvPT_,comdat
	.weak	_Z10testKernelIdEvPT_
	.p2align	3, 0x0
_Z10testKernelIdEvPT_:
	.quad	_Z25__device_stub__testKernelIdEvPT_
	.size	_Z10testKernelIdEvPT_, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z10testKernelIjEvPT_"
	.size	.L__unnamed_2, 22

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

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z10testKernelIfEvPT_"
	.size	.L__unnamed_4, 22

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z10testKernelIdEvPT_"
	.size	.L__unnamed_5, 22

	.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

	.type	.Lstr.30,@object                # @str.30
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr.30:
	.asciz	"atomicXor failed"
	.size	.Lstr.30, 17

	.type	.Lstr.31,@object                # @str.31
.Lstr.31:
	.asciz	"atomicOr failed"
	.size	.Lstr.31, 16

	.type	.Lstr.32,@object                # @str.32
.Lstr.32:
	.asciz	"atomicAnd failed"
	.size	.Lstr.32, 17

	.type	.Lstr.41,@object                # @str.41
.Lstr.41:
	.asciz	"atomicExch failed"
	.size	.Lstr.41, 18

	.type	.Lstr.42,@object                # @str.42
.Lstr.42:
	.asciz	"atomicCAS failed"
	.size	.Lstr.42, 17

	.type	.Lstr.43,@object                # @str.43
.Lstr.43:
	.asciz	"atomicDec failed"
	.size	.Lstr.43, 17

	.type	.Lstr.44,@object                # @str.44
.Lstr.44:
	.asciz	"atomicInc failed"
	.size	.Lstr.44, 17

	.type	.Lstr.45,@object                # @str.45
.Lstr.45:
	.asciz	"atomicMin failed"
	.size	.Lstr.45, 17

	.type	.Lstr.46,@object                # @str.46
.Lstr.46:
	.asciz	"atomicMax failed"
	.size	.Lstr.46, 17

	.type	.Lstr.47,@object                # @str.47
.Lstr.47:
	.asciz	"atomicSub failed"
	.size	.Lstr.47, 17

	.type	.Lstr.48,@object                # @str.48
.Lstr.48:
	.asciz	"atomicAdd failed"
	.size	.Lstr.48, 17

	.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__testKernelIiEvPT_
	.addrsig_sym _Z25__device_stub__testKernelIjEvPT_
	.addrsig_sym _Z25__device_stub__testKernelIyEvPT_
	.addrsig_sym _Z25__device_stub__testKernelIfEvPT_
	.addrsig_sym _Z25__device_stub__testKernelIdEvPT_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z10testKernelIiEvPT_
	.addrsig_sym _Z10testKernelIjEvPT_
	.addrsig_sym _Z10testKernelIyEvPT_
	.addrsig_sym _Z10testKernelIfEvPT_
	.addrsig_sym _Z10testKernelIdEvPT_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.section	.text._Z10testKernelIiEvPT_,"axG",@progbits,_Z10testKernelIiEvPT_,comdat
	.protected	_Z10testKernelIiEvPT_   ; -- Begin function _Z10testKernelIiEvPT_
	.globl	_Z10testKernelIiEvPT_
	.p2align	8
	.type	_Z10testKernelIiEvPT_,@function
_Z10testKernelIiEvPT_:                  ; @_Z10testKernelIiEvPT_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_mov_b32 s5, exec_lo
	s_mov_b32 s4, exec_lo
	v_mbcnt_lo_u32_b32 v1, s5, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_bcnt1_i32_b32 s5, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s5, s5, 10
	v_dual_mov_b32 v1, 0 :: v_dual_mov_b32 v2, s5
	s_waitcnt lgkmcnt(0)
	global_atomic_add_u32 v1, v2, s[2:3]
.LBB0_2:
	s_or_b32 exec_lo, exec_lo, s4
	s_load_b32 s0, s[0:1], 0x14
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s0, v[0:1]
	s_brev_b32 s0, 1
.LBB0_3:                                ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s1, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v1, s1
	s_lshl_b32 s1, 1, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s4, s4, s1
	s_max_i32 s0, s0, s5
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB0_3
; %bb.4:                                ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_6
; %bb.5:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s0
	global_atomic_max_i32 v0, v2, s[2:3] offset:12
.LBB0_6:
	s_or_b32 exec_lo, exec_lo, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s1, exec_lo
	s_brev_b32 s0, -2
.LBB0_7:                                ; %ComputeLoop4
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s4, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v1, s4
	s_lshl_b32 s4, 1, s4
	s_and_not1_b32 s1, s1, s4
	s_delay_alu instid0(VALU_DEP_1)
	s_min_i32 s0, s0, s5
	s_cmp_lg_u32 s1, 0
	s_cbranch_scc1 .LBB0_7
; %bb.8:                                ; %ComputeEnd5
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_10
; %bb.9:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s0
	global_atomic_min_i32 v0, v2, s[2:3] offset:16
.LBB0_10:
	s_or_b32 exec_lo, exec_lo, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, 17
	v_mov_b32_e32 v4, 0x89
	v_add_nc_u32_e32 v2, -1, v1
	s_mov_b32 s0, exec_lo
	s_clause 0x2
	global_atomic_inc_u32 v0, v3, s[2:3] offset:20
	global_atomic_dec_u32 v0, v4, s[2:3] offset:24
	global_atomic_cmpswap_b32 v0, v[1:2], s[2:3] offset:28
	v_lshl_add_u32 v0, v1, 1, 7
	s_mov_b32 s1, -1
.LBB0_11:                               ; %ComputeLoop8
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s4, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v0, s4
	s_lshl_b32 s4, 1, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s0, s0, s4
	s_and_b32 s1, s1, s5
	s_cmp_lg_u32 s0, 0
	s_cbranch_scc1 .LBB0_11
; %bb.12:                               ; %ComputeEnd9
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s0, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB0_14
; %bb.13:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s1
	global_atomic_and_b32 v0, v2, s[2:3] offset:32
.LBB0_14:
	s_or_b32 exec_lo, exec_lo, s4
	v_lshlrev_b32_e64 v0, v1, 1
	s_mov_b32 s1, exec_lo
.LBB0_15:                               ; %ComputeLoop12
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_ctz_i32_b32 s4, s1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v0, s4
	s_lshl_b32 s4, 1, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s1, s1, s4
	s_or_b32 s0, s0, s5
	s_cmp_lg_u32 s1, 0
	s_cbranch_scc1 .LBB0_15
; %bb.16:                               ; %ComputeEnd13
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s1, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB0_18
; %bb.17:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s0
	global_atomic_or_b32 v0, v2, s[2:3] offset:36
.LBB0_18:
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
.LBB0_19:                               ; %ComputeLoop16
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ctz_i32_b32 s4, s0
	v_readlane_b32 s5, v1, s4
	s_lshl_b32 s4, 1, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s0, s0, s4
	s_xor_b32 s1, s1, s5
	s_cmp_lg_u32 s0, 0
	s_cbranch_scc1 .LBB0_19
; %bb.20:                               ; %ComputeEnd17
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_22
; %bb.21:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s1
	global_atomic_xor_b32 v0, v2, s[2:3] offset:40
.LBB0_22:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_mov_b32 s1, exec_lo
	s_mov_b32 s0, exec_lo
	v_mbcnt_lo_u32_b32 v0, s1, 0
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB0_24
; %bb.23:
	s_bcnt1_i32_b32 s1, s1
	v_mov_b32_e32 v0, 0
	s_mul_i32 s1, s1, -10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v2, s1
	global_atomic_add_u32 v0, v2, s[2:3] offset:4
.LBB0_24:
	s_or_b32 exec_lo, exec_lo, s0
	v_mov_b32_e32 v0, 0
	global_atomic_swap_b32 v0, v1, s[2:3] offset:8
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelIiEvPT_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 264
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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
	.section	.text._Z10testKernelIiEvPT_,"axG",@progbits,_Z10testKernelIiEvPT_,comdat
.Lfunc_end0:
	.size	_Z10testKernelIiEvPT_, .Lfunc_end0-_Z10testKernelIiEvPT_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 744
; NumSgprs: 16
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z10testKernelIjEvPT_,"axG",@progbits,_Z10testKernelIjEvPT_,comdat
	.protected	_Z10testKernelIjEvPT_   ; -- Begin function _Z10testKernelIjEvPT_
	.globl	_Z10testKernelIjEvPT_
	.p2align	8
	.type	_Z10testKernelIjEvPT_,@function
_Z10testKernelIjEvPT_:                  ; @_Z10testKernelIjEvPT_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_mov_b32 s5, exec_lo
	s_mov_b32 s4, 0
	v_mbcnt_lo_u32_b32 v1, s5, 0
	s_mov_b32 s6, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB1_2
; %bb.1:
	s_bcnt1_i32_b32 s5, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s5, s5, 10
	v_dual_mov_b32 v1, 0 :: v_dual_mov_b32 v2, s5
	s_waitcnt lgkmcnt(0)
	global_atomic_add_u32 v1, v2, s[2:3]
.LBB1_2:
	s_or_b32 exec_lo, exec_lo, s6
	s_load_b32 s0, s[0:1], 0x14
	s_mov_b32 s5, exec_lo
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s0, v[0:1]
.LBB1_3:                                ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s0, s5
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s1, v1, s0
	s_lshl_b32 s0, 1, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s5, s5, s0
	s_max_u32 s4, s4, s1
	s_cmp_lg_u32 s5, 0
	s_cbranch_scc1 .LBB1_3
; %bb.4:                                ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB1_6
; %bb.5:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s4
	global_atomic_max_u32 v0, v2, s[2:3] offset:12
.LBB1_6:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s1, exec_lo
	s_mov_b32 s0, -1
.LBB1_7:                                ; %ComputeLoop4
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s4, s1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v1, s4
	s_lshl_b32 s4, 1, s4
	s_and_not1_b32 s1, s1, s4
	s_delay_alu instid0(VALU_DEP_1)
	s_min_u32 s0, s0, s5
	s_cmp_lg_u32 s1, 0
	s_cbranch_scc1 .LBB1_7
; %bb.8:                                ; %ComputeEnd5
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s1, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB1_10
; %bb.9:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s0
	global_atomic_min_u32 v0, v2, s[2:3] offset:16
.LBB1_10:
	s_or_b32 exec_lo, exec_lo, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, 17
	v_mov_b32_e32 v4, 0x89
	v_add_nc_u32_e32 v2, -1, v1
	s_mov_b32 s0, exec_lo
	s_clause 0x2
	global_atomic_inc_u32 v0, v3, s[2:3] offset:20
	global_atomic_dec_u32 v0, v4, s[2:3] offset:24
	global_atomic_cmpswap_b32 v0, v[1:2], s[2:3] offset:28
	v_lshl_add_u32 v0, v1, 1, 7
	s_mov_b32 s1, -1
.LBB1_11:                               ; %ComputeLoop8
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s4, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v0, s4
	s_lshl_b32 s4, 1, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s0, s0, s4
	s_and_b32 s1, s1, s5
	s_cmp_lg_u32 s0, 0
	s_cbranch_scc1 .LBB1_11
; %bb.12:                               ; %ComputeEnd9
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s0, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB1_14
; %bb.13:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s1
	global_atomic_and_b32 v0, v2, s[2:3] offset:32
.LBB1_14:
	s_or_b32 exec_lo, exec_lo, s4
	v_lshlrev_b32_e64 v0, v1, 1
	s_mov_b32 s1, exec_lo
.LBB1_15:                               ; %ComputeLoop12
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_ctz_i32_b32 s4, s1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s5, v0, s4
	s_lshl_b32 s4, 1, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s1, s1, s4
	s_or_b32 s0, s0, s5
	s_cmp_lg_u32 s1, 0
	s_cbranch_scc1 .LBB1_15
; %bb.16:                               ; %ComputeEnd13
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s1, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB1_18
; %bb.17:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s0
	global_atomic_or_b32 v0, v2, s[2:3] offset:36
.LBB1_18:
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s0, exec_lo
.LBB1_19:                               ; %ComputeLoop16
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ctz_i32_b32 s4, s0
	v_readlane_b32 s5, v1, s4
	s_lshl_b32 s4, 1, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s0, s0, s4
	s_xor_b32 s1, s1, s5
	s_cmp_lg_u32 s0, 0
	s_cbranch_scc1 .LBB1_19
; %bb.20:                               ; %ComputeEnd17
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB1_22
; %bb.21:
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v2, s1
	global_atomic_xor_b32 v0, v2, s[2:3] offset:40
.LBB1_22:
	s_or_b32 exec_lo, exec_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_mov_b32 s1, exec_lo
	s_mov_b32 s0, exec_lo
	v_mbcnt_lo_u32_b32 v0, s1, 0
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB1_24
; %bb.23:
	s_bcnt1_i32_b32 s1, s1
	v_mov_b32_e32 v0, 0
	s_mul_i32 s1, s1, -10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v2, s1
	global_atomic_add_u32 v0, v2, s[2:3] offset:4
.LBB1_24:
	s_or_b32 exec_lo, exec_lo, s0
	v_mov_b32_e32 v0, 0
	global_atomic_swap_b32 v0, v1, s[2:3] offset:8
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelIjEvPT_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 264
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 16
		.amdhsa_reserve_vcc 0
		.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
	.section	.text._Z10testKernelIjEvPT_,"axG",@progbits,_Z10testKernelIjEvPT_,comdat
.Lfunc_end1:
	.size	_Z10testKernelIjEvPT_, .Lfunc_end1-_Z10testKernelIjEvPT_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 744
; NumSgprs: 16
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 16
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z10testKernelIyEvPT_,"axG",@progbits,_Z10testKernelIyEvPT_,comdat
	.protected	_Z10testKernelIyEvPT_   ; -- Begin function _Z10testKernelIyEvPT_
	.globl	_Z10testKernelIyEvPT_
	.p2align	8
	.type	_Z10testKernelIyEvPT_,@function
_Z10testKernelIyEvPT_:                  ; @_Z10testKernelIyEvPT_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_mov_b32 s5, exec_lo
	s_mov_b32 s4, exec_lo
	v_mbcnt_lo_u32_b32 v1, s5, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB2_2
; %bb.1:
	s_bcnt1_i32_b32 s5, s5
	v_mov_b32_e32 v2, 0
	s_mul_i32 s5, s5, 10
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mov_b32_e32 v1, s5
	s_waitcnt lgkmcnt(0)
	global_atomic_add_u64 v2, v[1:2], s[2:3]
.LBB2_2:
	s_or_b32 exec_lo, exec_lo, s4
	s_load_b32 s0, s[0:1], 0x14
	s_mov_b32 s4, exec_lo
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s0, v[0:1]
	v_mov_b32_e32 v2, 0
	s_mov_b64 s[0:1], 0
.LBB2_3:                                ; %ComputeLoop
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s5, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s7, v2, s5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_readlane_b32 s6, v1, s5
	v_cmp_gt_u64_e64 s8, s[0:1], s[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_and_b32 s8, s8, exec_lo
	s_cselect_b32 s1, s1, s7
	s_cselect_b32 s0, s0, s6
	s_lshl_b32 s5, 1, s5
	s_and_not1_b32 s4, s4, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB2_3
; %bb.4:                                ; %ComputeEnd
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB2_6
; %bb.5:
	v_mov_b32_e32 v4, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s0
	global_atomic_max_u64 v0, v[3:4], s[2:3] offset:24
.LBB2_6:
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s4, exec_lo
	s_mov_b64 s[0:1], -1
.LBB2_7:                                ; %ComputeLoop4
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s5, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_readlane_b32 s7, v2, s5
	v_readlane_b32 s6, v1, s5
	v_cmp_lt_u64_e64 s8, s[0:1], s[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	s_and_b32 s8, s8, exec_lo
	s_cselect_b32 s1, s1, s7
	s_cselect_b32 s0, s0, s6
	s_lshl_b32 s5, 1, s5
	s_and_not1_b32 s4, s4, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB2_7
; %bb.8:                                ; %ComputeEnd5
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB2_10
; %bb.9:
	v_mov_b32_e32 v4, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s0
	global_atomic_min_u64 v0, v[3:4], s[2:3] offset:32
.LBB2_10:
	s_or_b32 exec_lo, exec_lo, s4
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v5, 17
	v_add_co_u32 v3, s0, v1, -1
	v_mov_b32_e32 v6, 0x89
	v_add_co_ci_u32_e64 v4, null, 0, -1, s0
	s_clause 0x2
	global_atomic_inc_u32 v0, v5, s[2:3] offset:40
	global_atomic_dec_u32 v0, v6, s[2:3] offset:48
	global_atomic_cmpswap_b64 v0, v[1:4], s[2:3] offset:56
	v_lshlrev_b64 v[3:4], 1, v[1:2]
	s_mov_b32 s4, exec_lo
	s_mov_b64 s[0:1], -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, v3, 7
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v4, vcc_lo
.LBB2_11:                               ; %ComputeLoop8
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s5, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s7, v3, s5
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s6, v0, s5
	s_lshl_b32 s5, 1, s5
	s_and_not1_b32 s4, s4, s5
	s_delay_alu instid0(VALU_DEP_1)
	s_and_b64 s[0:1], s[0:1], s[6:7]
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB2_11
; %bb.12:                               ; %ComputeEnd9
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB2_14
; %bb.13:
	v_mov_b32_e32 v4, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s0
	global_atomic_and_b64 v0, v[3:4], s[2:3] offset:64
.LBB2_14:
	s_or_b32 exec_lo, exec_lo, s4
	v_lshlrev_b32_e64 v0, v1, 1
	s_mov_b32 s4, exec_lo
	s_mov_b64 s[0:1], 0
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v0
.LBB2_15:                               ; %ComputeLoop12
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s5, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s7, v3, s5
	v_readlane_b32 s6, v0, s5
	s_lshl_b32 s5, 1, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_not1_b32 s4, s4, s5
	s_or_b64 s[0:1], s[0:1], s[6:7]
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB2_15
; %bb.16:                               ; %ComputeEnd13
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB2_18
; %bb.17:
	v_mov_b32_e32 v4, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s0
	global_atomic_or_b64 v0, v[3:4], s[2:3] offset:72
.LBB2_18:
	s_or_b32 exec_lo, exec_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mov_b32 s4, exec_lo
	s_mov_b64 s[0:1], 0
.LBB2_19:                               ; %ComputeLoop16
                                        ; =>This Inner Loop Header: Depth=1
	s_ctz_i32_b32 s5, s4
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s7, v2, s5
	v_readlane_b32 s6, v1, s5
	s_lshl_b32 s5, 1, s5
	s_and_not1_b32 s4, s4, s5
	s_delay_alu instid0(VALU_DEP_1)
	s_xor_b64 s[0:1], s[0:1], s[6:7]
	s_cmp_lg_u32 s4, 0
	s_cbranch_scc1 .LBB2_19
; %bb.20:                               ; %ComputeEnd17
	v_mbcnt_lo_u32_b32 v0, exec_lo, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_xor_b32 s4, exec_lo, s4
	s_cbranch_execz .LBB2_22
; %bb.21:
	v_mov_b32_e32 v4, s1
	v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v3, s0
	global_atomic_xor_b64 v0, v[3:4], s[2:3] offset:80
.LBB2_22:
	s_or_b32 exec_lo, exec_lo, s4
	v_mov_b32_e32 v0, 0
	global_atomic_swap_b64 v0, v[1:2], s[2:3] offset:16
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelIyEvPT_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 264
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 7
		.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
	.section	.text._Z10testKernelIyEvPT_,"axG",@progbits,_Z10testKernelIyEvPT_,comdat
.Lfunc_end2:
	.size	_Z10testKernelIyEvPT_, .Lfunc_end2-_Z10testKernelIyEvPT_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 828
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z10testKernelIfEvPT_,"axG",@progbits,_Z10testKernelIfEvPT_,comdat
	.protected	_Z10testKernelIfEvPT_   ; -- Begin function _Z10testKernelIfEvPT_
	.globl	_Z10testKernelIfEvPT_
	.p2align	8
	.type	_Z10testKernelIfEvPT_,@function
_Z10testKernelIfEvPT_:                  ; @_Z10testKernelIfEvPT_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x0
	s_mov_b32 s6, exec_lo
	s_mov_b32 s5, 0
	v_mbcnt_lo_u32_b32 v1, s6, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v1
	s_cbranch_execz .LBB3_3
; %bb.1:
	s_waitcnt lgkmcnt(0)
	s_load_b32 s7, s[2:3], 0x0
	s_bcnt1_i32_b32 s6, s6
	v_mov_b32_e32 v4, 0
	v_cvt_f32_ubyte0_e32 v1, s6
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mul_f32 v3, 0x41200000, v1 :: v_dual_mov_b32 v2, s7
.LBB3_2:                                ; %atomicrmw.start
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v1, v2, v3
	global_atomic_cmpswap_b32 v1, v4, v[1:2], s[2:3] glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, v1, v2
	v_mov_b32_e32 v2, v1
	s_or_b32 s5, vcc_lo, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s5
	s_cbranch_execnz .LBB3_2
.LBB3_3:                                ; %Flow5
	s_or_b32 exec_lo, exec_lo, s4
	s_load_b32 s0, s[0:1], 0x14
	s_waitcnt lgkmcnt(0)
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s0, v[0:1]
	v_cvt_f32_u32_e32 v0, v1
	v_mov_b32_e32 v1, 0
	global_atomic_swap_b32 v1, v0, s[2:3] offset:8
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelIfEvPT_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 264
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.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
	.section	.text._Z10testKernelIfEvPT_,"axG",@progbits,_Z10testKernelIfEvPT_,comdat
.Lfunc_end3:
	.size	_Z10testKernelIfEvPT_, .Lfunc_end3-_Z10testKernelIfEvPT_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 192
; NumSgprs: 18
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 5
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.section	.text._Z10testKernelIdEvPT_,"axG",@progbits,_Z10testKernelIdEvPT_,comdat
	.protected	_Z10testKernelIdEvPT_   ; -- Begin function _Z10testKernelIdEvPT_
	.globl	_Z10testKernelIdEvPT_
	.p2align	8
	.type	_Z10testKernelIdEvPT_,@function
_Z10testKernelIdEvPT_:                  ; @_Z10testKernelIdEvPT_
; %bb.0:
	s_mov_b32 s3, exec_lo
	s_mov_b32 s2, 0
	v_mbcnt_lo_u32_b32 v0, s3, 0
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB4_3
; %bb.1:
	s_bcnt1_i32_b32 s3, s3
	s_load_b64 s[0:1], s[0:1], 0x0
	v_cvt_f64_u32_e32 v[0:1], s3
	v_mov_b32_e32 v6, 0
	s_waitcnt lgkmcnt(0)
	s_load_b64 s[4:5], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_2)
	v_mul_f64 v[4:5], 0x40240000, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_dual_mov_b32 v2, s4 :: v_dual_mov_b32 v3, s5
.LBB4_2:                                ; %atomicrmw.start
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[0:1], v[2:3], v[4:5]
	global_atomic_cmpswap_b64 v[0:1], v6, v[0:3], s[0:1] glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u64_e32 vcc_lo, v[0:1], v[2:3]
	v_dual_mov_b32 v3, v1 :: v_dual_mov_b32 v2, v0
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB4_2
.LBB4_3:
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10testKernelIdEvPT_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 8
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 7
		.amdhsa_next_free_sgpr 6
		.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
	.section	.text._Z10testKernelIdEvPT_,"axG",@progbits,_Z10testKernelIdEvPT_,comdat
.Lfunc_end4:
	.size	_Z10testKernelIdEvPT_, .Lfunc_end4-_Z10testKernelIdEvPT_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 148
; NumSgprs: 8
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 8
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.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
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         12
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         20
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         22
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         24
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         26
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         28
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         30
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         72
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 264
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelIiEvPT_
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelIiEvPT_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         12
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         20
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         22
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         24
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         26
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         28
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         30
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         72
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 264
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelIjEvPT_
    .private_segment_fixed_size: 0
    .sgpr_count:     16
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelIjEvPT_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         12
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         20
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         22
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         24
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         26
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         28
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         30
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         72
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 264
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelIyEvPT_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelIyEvPT_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         12
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         20
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         22
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         24
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         26
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         28
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         30
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         72
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 264
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelIfEvPT_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelIfEvPT_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 8
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10testKernelIdEvPT_
    .private_segment_fixed_size: 0
    .sgpr_count:     8
    .sgpr_spill_count: 0
    .symbol:         _Z10testKernelIdEvPT_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
