// Template file for the OpenCL Assignment 4

#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>
#include <math.h>
#include <sys/time.h>


#define NUM_PARTICLES 10000
#define NUM_ITERATIONS 1000
#define BLOCK_SIZE 64
#define error 1e-6
// This is a macro for checking the error variable:>.
#define CHK_ERROR(err) if (err != CL_SUCCESS) fprintf(stderr,"Error: %s\n",clGetErrorString(err));

// A errorCode to string converter (forward declaration)
const char* clGetErrorString(int);




const char *mykernel = 
"typedef struct{                                         \n"
    "float position[3];                                    \n"
    "float velocity[3];                                   \n"
"}particle;                                                       \n"
"__kernel                                                 \n" 
"void kernel_upload(                                       \n"
"   const int iter,                                          \n"
"   const int nump,                                        \n"    
"   __global particle* X)                                    \n"                                      
"{int index = get_global_id(0);                           \n"
"if(index < nump) {                                                \n"
"   X[index].velocity[0] = (3*index + iter) % nump;             \n"
"   X[index].velocity[1] = (4*index + iter) % nump;            \n"
"   X[index].velocity[2] = (5*index + iter) % nump;             \n"
"   X[index].position[0] = X[index].position[0] + X[index].velocity[0];             \n"    
"   X[index].position[1] = X[index].position[1] + X[index].velocity[1];             \n"
"   X[index].position[2] = X[index].position[2] + X[index].velocity[2];             \n"
"}                                                          \n"
"}                                                        \n";

typedef struct{
    float position[3];
    float velocity[3];
}particle;

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp,NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
}

void init_Array(particle *particulas){
    for(int i = 0; i < NUM_PARTICLES; i++){
        particulas[i].position[0] = rand() % 1000;
        particulas[i].position[1] = rand() % 1000;
        particulas[i].position[2] = rand() % 1000;

        particulas[i].velocity[0] = rand() % 1000;
        particulas[i].velocity[1] = rand() % 1000;
        particulas[i].velocity[2] = rand() % 1000;
    }
}


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

    //All the variables used
    particle X[NUM_PARTICLES],  results[NUM_PARTICLES];
    float array = NUM_PARTICLES * sizeof(particle);
    int count = NUM_PARTICLES;
    bool correct = true;
    int i, j;

    cl_platform_id * platforms; cl_uint     n_platform;

    // Find OpenCL Platforms
    cl_int err = clGetPlatformIDs(0, NULL, &n_platform); CHK_ERROR(err);
    platforms = (cl_platform_id *) malloc(sizeof(cl_platform_id)*n_platform);
    err = clGetPlatformIDs(n_platform, platforms, NULL); CHK_ERROR(err);

    // Find and sort devices
    cl_device_id *device_list; cl_uint n_devices;
    err = clGetDeviceIDs( platforms[0], CL_DEVICE_TYPE_GPU, 0,NULL, &n_devices);CHK_ERROR(err);
    device_list = (cl_device_id *) malloc(sizeof(cl_device_id)*n_devices);
    err = clGetDeviceIDs( platforms[0],CL_DEVICE_TYPE_GPU, n_devices, device_list, NULL);CHK_ERROR(err);
    
    // Create and initialize an OpenCL context
    cl_context context = clCreateContext( NULL, n_devices, device_list, NULL, NULL, &err);CHK_ERROR(err);

    if (!context) {
        printf("Error: Failed to create a compute context!\n");
        return EXIT_FAILURE;
    }

    // Create a command queue
    cl_command_queue cmd_queue = clCreateCommandQueue(context, device_list[0], 0, &err);CHK_ERROR(err);

    if (!cmd_queue) {
        printf("Error: Failed to create a command queue!\n");
        return EXIT_FAILURE;
    }


    /* Insert your own code here */

    //program
    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&mykernel, NULL, &err);CHK_ERROR(err);


    //Initialize values of the arrays
 
    init_Array(X);

    //Array on GPU part
    cl_mem dx = clCreateBuffer(context, CL_MEM_READ_WRITE, array, NULL, &err);CHK_ERROR(err);


    //Copy the values of X in dx
    err = clEnqueueWriteBuffer(cmd_queue, dx, CL_TRUE, 0, array, X, 0, NULL, NULL);CHK_ERROR(err);


    //Cpu part
    double start = cpuSecond();
    for(j = 0; j < NUM_ITERATIONS; j++) {
        for(i = 0; i < NUM_PARTICLES; i++){
            X[i].velocity[0] = (3*i + j) % NUM_PARTICLES;             
            X[i].velocity[1] = (4*i + j) % NUM_PARTICLES;            
            X[i].velocity[2] = (5*i + j) % NUM_PARTICLES;             
            X[i].position[0] = X[i].position[0] + X[i].velocity[0];                 
            X[i].position[1] = X[i].position[1] + X[i].velocity[1];            
            X[i].position[2] = X[i].position[2] + X[i].velocity[2]; 
        }
    }
    double totalCpu = cpuSecond() - start;        

    printf("Computing SAXPY on the CPU... Done\n");
    
    //build the program 
    err = clBuildProgram(program, 1, device_list, NULL, NULL, NULL);
    if (err != CL_SUCCESS){
        size_t len;
        char buffer[2048];
        clGetProgramBuildInfo(program, device_list[0], CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        fprintf(stderr, "Build error: %s\n", buffer);
        return 0;
    }

    //Create the kernel function
    cl_kernel kernel = clCreateKernel(program, "kernel_upload", &err);CHK_ERROR(err);
    


    //More argument, number of workitems we are going to use
    unsigned int mult = BLOCK_SIZE;
    unsigned int size = BLOCK_SIZE;
    while(mult < NUM_PARTICLES){

        mult = mult * 2;
    }
    size_t n_workitem[1] = {mult};
    size_t workgroup_size[1] = {size};

    err = clSetKernelArg(kernel, 1, sizeof(int), (void*) &count);CHK_ERROR(err);
    err = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &dx);CHK_ERROR(err);

    
    double st = cpuSecond();
    for(i = 0; i < NUM_ITERATIONS; i++){

        //Arguments for the Kernel function
        err = clSetKernelArg(kernel, 0, sizeof(int), (void*) &i);CHK_ERROR(err);



        //Launch kernel
        err = clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL, n_workitem, workgroup_size, 0, NULL, NULL);CHK_ERROR(err);


        //Wait for everything to finish
        err = clFlush(cmd_queue);
        err = clFinish(cmd_queue);
    }
    err = clEnqueueReadBuffer(cmd_queue, dx, CL_TRUE, 0, array, results, 0, NULL, NULL);CHK_ERROR(err);
    double timeGpu = cpuSecond() - st;
    printf("Computing SACY on the GPU... Done\n");

    for(int i = 0; i < NUM_PARTICLES; i++){
        for(int dim = 0; dim < 3; dim++){
            if(abs(results[i].position[dim] - X[i].position[dim]) > error ){
                correct = false;
                printf("Error en %d i en dim %d, resultado %f i %f\n",i,dim,results[i].position[1],X[i].position[1]);
                break;
            }
        }
        if(!correct) {break;}
    }

    if (correct) {
        printf("Comparing the output for each implementation... Correct!\n");
    }
    else {
        printf("Comparing the output for each implementation... Incorrect!\n");
    }

    printf("NUM_PARTICLES: %d\n", NUM_PARTICLES);
    printf("BLOCK_SIZE: %d\n", BLOCK_SIZE);
    printf("NUM_ITERATIONS: %d\n", NUM_ITERATIONS);
    printf("Execution time on CPU: %f\n", totalCpu);
    printf("Execution time on GPU: %f\n", timeGpu);


    // Finally, release all that we have allocated.
    clReleaseMemObject(dx);CHK_ERROR(err);
    clReleaseProgram(program);CHK_ERROR(err);
    clReleaseKernel(kernel);CHK_ERROR(err);
    err = clReleaseCommandQueue(cmd_queue);CHK_ERROR(err);
    err = clReleaseContext(context);CHK_ERROR(err);
    free(platforms);
    free(device_list);
    
    return 0;
}



// The source for this particular version is from: https://stackoverflow.com/questions/24326432/convenient-way-to-show-opencl-error-codes
const char* clGetErrorString(int errorCode) {
    switch (errorCode) {
    case 0: return "CL_SUCCESS";
    case -1: return "CL_DEVICE_NOT_FOUND";
    case -2: return "CL_DEVICE_NOT_AVAILABLE";
    case -3: return "CL_COMPILER_NOT_AVAILABLE";
    case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
    case -5: return "CL_OUT_OF_RESOURCES";
    case -6: return "CL_OUT_OF_HOST_MEMORY";
    case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE";
    case -8: return "CL_MEM_COPY_OVERLAP";
    case -9: return "CL_IMAGE_FORMAT_MISMATCH";
    case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
    case -12: return "CL_MAP_FAILURE";
    case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
    case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
    case -15: return "CL_COMPILE_PROGRAM_FAILURE";
    case -16: return "CL_LINKER_NOT_AVAILABLE";
    case -17: return "CL_LINK_PROGRAM_FAILURE";
    case -18: return "CL_DEVICE_PARTITION_FAILED";
    case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
    case -30: return "CL_INVALID_VALUE";
    case -31: return "CL_INVALID_DEVICE_TYPE";
    case -32: return "CL_INVALID_PLATFORM";
    case -33: return "CL_INVALID_DEVICE";
    case -34: return "CL_INVALID_CONTEXT";
    case -35: return "CL_INVALID_QUEUE_PROPERTIES";
    case -36: return "CL_INVALID_COMMAND_QUEUE";
    case -37: return "CL_INVALID_HOST_PTR";
    case -38: return "CL_INVALID_MEM_OBJECT";
    case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
    case -40: return "CL_INVALID_IMAGE_SIZE";
    case -41: return "CL_INVALID_SAMPLER";
    case -42: return "CL_INVALID_BINARY";
    case -43: return "CL_INVALID_BUILD_OPTIONS";
    case -44: return "CL_INVALID_PROGRAM";
    case -45: return "CL_INVALID_PROGRAM_EXECUTABLE";
    case -46: return "CL_INVALID_KERNEL_NAME";
    case -47: return "CL_INVALID_KERNEL_DEFINITION";
    case -48: return "CL_INVALID_KERNEL";
    case -49: return "CL_INVALID_ARG_INDEX";
    case -50: return "CL_INVALID_ARG_VALUE";
    case -51: return "CL_INVALID_ARG_SIZE";
    case -52: return "CL_INVALID_KERNEL_ARGS";
    case -53: return "CL_INVALID_WORK_DIMENSION";
    case -54: return "CL_INVALID_WORK_GROUP_SIZE";
    case -55: return "CL_INVALID_WORK_ITEM_SIZE";
    case -56: return "CL_INVALID_GLOBAL_OFFSET";
    case -57: return "CL_INVALID_EVENT_WAIT_LIST";
    case -58: return "CL_INVALID_EVENT";
    case -59: return "CL_INVALID_OPERATION";
    case -60: return "CL_INVALID_GL_OBJECT";
    case -61: return "CL_INVALID_BUFFER_SIZE";
    case -62: return "CL_INVALID_MIP_LEVEL";
    case -63: return "CL_INVALID_GLOBAL_WORK_SIZE";
    case -64: return "CL_INVALID_PROPERTY";
    case -65: return "CL_INVALID_IMAGE_DESCRIPTOR";
    case -66: return "CL_INVALID_COMPILER_OPTIONS";
    case -67: return "CL_INVALID_LINKER_OPTIONS";
    case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT";
    case -69: return "CL_INVALID_PIPE_SIZE";
    case -70: return "CL_INVALID_DEVICE_QUEUE";
    case -71: return "CL_INVALID_SPEC_ID";
    case -72: return "CL_MAX_SIZE_RESTRICTION_EXCEEDED";
    case -1002: return "CL_INVALID_D3D10_DEVICE_KHR";
    case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR";
    case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
    case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
    case -1006: return "CL_INVALID_D3D11_DEVICE_KHR";
    case -1007: return "CL_INVALID_D3D11_RESOURCE_KHR";
    case -1008: return "CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR";
    case -1009: return "CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR";
    case -1010: return "CL_INVALID_DX9_MEDIA_ADAPTER_KHR";
    case -1011: return "CL_INVALID_DX9_MEDIA_SURFACE_KHR";
    case -1012: return "CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR";
    case -1013: return "CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR";
    case -1093: return "CL_INVALID_EGL_OBJECT_KHR";
    case -1092: return "CL_EGL_RESOURCE_NOT_ACQUIRED_KHR";
    case -1001: return "CL_PLATFORM_NOT_FOUND_KHR";
    case -1057: return "CL_DEVICE_PARTITION_FAILED_EXT";
    case -1058: return "CL_INVALID_PARTITION_COUNT_EXT";
    case -1059: return "CL_INVALID_PARTITION_NAME_EXT";
    case -1094: return "CL_INVALID_ACCELERATOR_INTEL";
    case -1095: return "CL_INVALID_ACCELERATOR_TYPE_INTEL";
    case -1096: return "CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL";
    case -1097: return "CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL";
    case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
    case -1098: return "CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL";
    case -1099: return "CL_INVALID_VA_API_MEDIA_SURFACE_INTEL";
    case -1100: return "CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL";
    case -1101: return "CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL";
    default: return "CL_UNKNOWN_ERROR";
    }
}


	.file	"tmpxft_002cbad8_00000000-6_exercise_3.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2554:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2554:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9cpuSecondv
	.type	_Z9cpuSecondv, @function
_Z9cpuSecondv:
.LFB2548:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	xorl	%esi, %esi
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	cvtsi2sdq	16(%rsp), %xmm0
	mulsd	.LC0(%rip), %xmm0
	cvtsi2sdq	8(%rsp), %xmm1
	addsd	%xmm1, %xmm0
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2548:
	.size	_Z9cpuSecondv, .-_Z9cpuSecondv
	.globl	_Z10init_ArrayP8particle
	.type	_Z10init_ArrayP8particle, @function
_Z10init_ArrayP8particle:
.LFB2549:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	leaq	240000(%rdi), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	$1000, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
.L7:
	call	rand@PLT
	addq	$24, %rbx
	cltd
	idivl	%ebp
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -24(%rbx)
	call	rand@PLT
	cltd
	idivl	%ebp
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -20(%rbx)
	call	rand@PLT
	cltd
	idivl	%ebp
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -16(%rbx)
	call	rand@PLT
	cltd
	idivl	%ebp
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -12(%rbx)
	call	rand@PLT
	cltd
	idivl	%ebp
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -8(%rbx)
	call	rand@PLT
	cltd
	idivl	%ebp
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -4(%rbx)
	cmpq	%r12, %rbx
	jne	.L7
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2549:
	.size	_Z10init_ArrayP8particle, .-_Z10init_ArrayP8particle
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"CL_SUCCESS"
.LC2:
	.string	"CL_UNKNOWN_ERROR"
.LC3:
	.string	"CL_DEVICE_NOT_AVAILABLE"
.LC4:
	.string	"CL_COMPILER_NOT_AVAILABLE"
.LC5:
	.string	"CL_MEM_OBJECT_ALLOCATION_FAILURE"
.LC6:
	.string	"CL_OUT_OF_RESOURCES"
.LC7:
	.string	"CL_OUT_OF_HOST_MEMORY"
.LC8:
	.string	"CL_PROFILING_INFO_NOT_AVAILABLE"
.LC9:
	.string	"CL_MEM_COPY_OVERLAP"
.LC10:
	.string	"CL_IMAGE_FORMAT_MISMATCH"
.LC11:
	.string	"CL_IMAGE_FORMAT_NOT_SUPPORTED"
.LC12:
	.string	"CL_MAP_FAILURE"
.LC13:
	.string	"CL_MISALIGNED_SUB_BUFFER_OFFSET"
.LC14:
	.string	"CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"
.LC15:
	.string	"CL_COMPILE_PROGRAM_FAILURE"
.LC16:
	.string	"CL_LINKER_NOT_AVAILABLE"
.LC17:
	.string	"CL_LINK_PROGRAM_FAILURE"
.LC18:
	.string	"CL_DEVICE_PARTITION_FAILED"
.LC19:
	.string	"CL_KERNEL_ARG_INFO_NOT_AVAILABLE"
.LC20:
	.string	"CL_INVALID_VALUE"
.LC21:
	.string	"CL_INVALID_DEVICE_TYPE"
.LC22:
	.string	"CL_INVALID_PLATFORM"
.LC23:
	.string	"CL_INVALID_DEVICE"
.LC24:
	.string	"CL_INVALID_CONTEXT"
.LC25:
	.string	"CL_INVALID_QUEUE_PROPERTIES"
.LC26:
	.string	"CL_INVALID_COMMAND_QUEUE"
.LC27:
	.string	"CL_INVALID_HOST_PTR"
.LC28:
	.string	"CL_INVALID_MEM_OBJECT"
.LC29:
	.string	"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"
.LC30:
	.string	"CL_INVALID_IMAGE_SIZE"
.LC31:
	.string	"CL_INVALID_SAMPLER"
.LC32:
	.string	"CL_INVALID_BINARY"
.LC33:
	.string	"CL_INVALID_BUILD_OPTIONS"
.LC34:
	.string	"CL_INVALID_PROGRAM"
.LC35:
	.string	"CL_INVALID_PROGRAM_EXECUTABLE"
.LC36:
	.string	"CL_INVALID_KERNEL_NAME"
.LC37:
	.string	"CL_INVALID_KERNEL_DEFINITION"
.LC38:
	.string	"CL_INVALID_KERNEL"
.LC39:
	.string	"CL_INVALID_ARG_INDEX"
.LC40:
	.string	"CL_INVALID_ARG_VALUE"
.LC41:
	.string	"CL_INVALID_ARG_SIZE"
.LC42:
	.string	"CL_INVALID_KERNEL_ARGS"
.LC43:
	.string	"CL_INVALID_WORK_DIMENSION"
.LC44:
	.string	"CL_INVALID_WORK_GROUP_SIZE"
.LC45:
	.string	"CL_INVALID_WORK_ITEM_SIZE"
.LC46:
	.string	"CL_INVALID_GLOBAL_OFFSET"
.LC47:
	.string	"CL_INVALID_EVENT_WAIT_LIST"
.LC48:
	.string	"CL_INVALID_EVENT"
.LC49:
	.string	"CL_INVALID_OPERATION"
.LC50:
	.string	"CL_INVALID_GL_OBJECT"
.LC51:
	.string	"CL_INVALID_BUFFER_SIZE"
.LC52:
	.string	"CL_INVALID_MIP_LEVEL"
.LC53:
	.string	"CL_INVALID_GLOBAL_WORK_SIZE"
.LC54:
	.string	"CL_INVALID_PROPERTY"
.LC55:
	.string	"CL_INVALID_IMAGE_DESCRIPTOR"
.LC56:
	.string	"CL_INVALID_COMPILER_OPTIONS"
.LC57:
	.string	"CL_INVALID_LINKER_OPTIONS"
.LC58:
	.string	"CL_INVALID_DEVICE_PARTITION_COUNT"
.LC59:
	.string	"CL_INVALID_PIPE_SIZE"
.LC60:
	.string	"CL_INVALID_DEVICE_QUEUE"
.LC61:
	.string	"CL_INVALID_SPEC_ID"
.LC62:
	.string	"CL_MAX_SIZE_RESTRICTION_EXCEEDED"
.LC63:
	.string	"CL_INVALID_D3D10_DEVICE_KHR"
.LC64:
	.string	"CL_INVALID_D3D10_RESOURCE_KHR"
.LC65:
	.string	"CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR"
.LC66:
	.string	"CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR"
.LC67:
	.string	"CL_INVALID_D3D11_DEVICE_KHR"
.LC68:
	.string	"CL_INVALID_D3D11_RESOURCE_KHR"
.LC69:
	.string	"CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR"
.LC70:
	.string	"CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR"
.LC71:
	.string	"CL_INVALID_DX9_MEDIA_ADAPTER_KHR"
.LC72:
	.string	"CL_INVALID_DX9_MEDIA_SURFACE_KHR"
.LC73:
	.string	"CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR"
.LC74:
	.string	"CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR"
.LC75:
	.string	"CL_INVALID_EGL_OBJECT_KHR"
.LC76:
	.string	"CL_EGL_RESOURCE_NOT_ACQUIRED_KHR"
.LC77:
	.string	"CL_PLATFORM_NOT_FOUND_KHR"
.LC78:
	.string	"CL_DEVICE_PARTITION_FAILED_EXT"
.LC79:
	.string	"CL_INVALID_PARTITION_COUNT_EXT"
.LC80:
	.string	"CL_INVALID_PARTITION_NAME_EXT"
.LC81:
	.string	"CL_INVALID_ACCELERATOR_INTEL"
.LC82:
	.string	"CL_INVALID_ACCELERATOR_TYPE_INTEL"
.LC83:
	.string	"CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL"
.LC84:
	.string	"CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL"
.LC85:
	.string	"CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR"
.LC86:
	.string	"CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL"
.LC87:
	.string	"CL_INVALID_VA_API_MEDIA_SURFACE_INTEL"
.LC88:
	.string	"CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL"
.LC89:
	.string	"CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL"
.LC90:
	.string	"CL_DEVICE_NOT_FOUND"
	.text
	.globl	_Z16clGetErrorStringi
	.type	_Z16clGetErrorStringi, @function
_Z16clGetErrorStringi:
.LFB2551:
	.cfi_startproc
	endbr64
	leaq	.LC2(%rip), %rax
	testl	%edi, %edi
	jg	.L10
	cmpl	$-72, %edi
	jge	.L12
	cmpl	$-999, %edi
	jge	.L10
	cmpl	$-1013, %edi
	jge	.L13
	cmpl	$-1091, %edi
	jge	.L14
	cmpl	$-1101, %edi
	jge	.L15
	ret
.L12:
	addl	$72, %edi
	cmpl	$72, %edi
	ja	.L16
	leaq	.L18(%rip), %rdx
	movslq	(%rdx,%rdi,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L18:
	.long	.L78-.L18
	.long	.L77-.L18
	.long	.L76-.L18
	.long	.L75-.L18
	.long	.L74-.L18
	.long	.L73-.L18
	.long	.L72-.L18
	.long	.L71-.L18
	.long	.L70-.L18
	.long	.L69-.L18
	.long	.L68-.L18
	.long	.L67-.L18
	.long	.L66-.L18
	.long	.L65-.L18
	.long	.L64-.L18
	.long	.L63-.L18
	.long	.L62-.L18
	.long	.L61-.L18
	.long	.L60-.L18
	.long	.L59-.L18
	.long	.L58-.L18
	.long	.L57-.L18
	.long	.L56-.L18
	.long	.L55-.L18
	.long	.L54-.L18
	.long	.L53-.L18
	.long	.L52-.L18
	.long	.L51-.L18
	.long	.L50-.L18
	.long	.L49-.L18
	.long	.L48-.L18
	.long	.L47-.L18
	.long	.L46-.L18
	.long	.L45-.L18
	.long	.L44-.L18
	.long	.L43-.L18
	.long	.L42-.L18
	.long	.L41-.L18
	.long	.L40-.L18
	.long	.L39-.L18
	.long	.L38-.L18
	.long	.L37-.L18
	.long	.L36-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L16-.L18
	.long	.L35-.L18
	.long	.L34-.L18
	.long	.L33-.L18
	.long	.L32-.L18
	.long	.L31-.L18
	.long	.L30-.L18
	.long	.L29-.L18
	.long	.L28-.L18
	.long	.L16-.L18
	.long	.L27-.L18
	.long	.L26-.L18
	.long	.L25-.L18
	.long	.L24-.L18
	.long	.L23-.L18
	.long	.L22-.L18
	.long	.L21-.L18
	.long	.L20-.L18
	.long	.L19-.L18
	.long	.L105-.L18
	.long	.L17-.L18
	.text
.L17:
	leaq	.LC1(%rip), %rax
	ret
.L16:
	leaq	.LC2(%rip), %rax
	ret
.L13:
	addl	$1012, %edi
	leaq	.LC74(%rip), %rax
	cmpl	$12, %edi
	ja	.L10
	leaq	.L81(%rip), %rdx
	movslq	(%rdx,%rdi,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L81:
	.long	.L92-.L81
	.long	.L91-.L81
	.long	.L90-.L81
	.long	.L89-.L81
	.long	.L88-.L81
	.long	.L87-.L81
	.long	.L86-.L81
	.long	.L85-.L81
	.long	.L84-.L81
	.long	.L83-.L81
	.long	.L106-.L81
	.long	.L82-.L81
	.long	.L80-.L81
	.text
.L83:
	leaq	.LC64(%rip), %rax
	ret
.L15:
	addl	$1100, %edi
	leaq	.LC89(%rip), %rax
	cmpl	$8, %edi
	ja	.L10
	leaq	.L95(%rip), %rdx
	movslq	(%rdx,%rdi,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L95:
	.long	.L102-.L95
	.long	.L101-.L95
	.long	.L100-.L95
	.long	.L99-.L95
	.long	.L98-.L95
	.long	.L97-.L95
	.long	.L96-.L95
	.long	.L107-.L95
	.long	.L94-.L95
	.text
.L94:
	leaq	.LC76(%rip), %rax
	ret
.L14:
	leaq	.LC79(%rip), %rax
	cmpl	$-1058, %edi
	je	.L10
	leaq	.LC78(%rip), %rax
	cmpl	$-1057, %edi
	je	.L10
	cmpl	$-1059, %edi
	leaq	.LC80(%rip), %rdx
	leaq	.LC2(%rip), %rax
	cmove	%rdx, %rax
	ret
.L19:
	leaq	.LC3(%rip), %rax
	ret
.L20:
	leaq	.LC4(%rip), %rax
	ret
.L21:
	leaq	.LC5(%rip), %rax
	ret
.L22:
	leaq	.LC6(%rip), %rax
	ret
.L23:
	leaq	.LC7(%rip), %rax
	ret
.L24:
	leaq	.LC8(%rip), %rax
	ret
.L25:
	leaq	.LC9(%rip), %rax
	ret
.L26:
	leaq	.LC10(%rip), %rax
	ret
.L27:
	leaq	.LC11(%rip), %rax
	ret
.L28:
	leaq	.LC12(%rip), %rax
	ret
.L29:
	leaq	.LC13(%rip), %rax
	ret
.L30:
	leaq	.LC14(%rip), %rax
	ret
.L31:
	leaq	.LC15(%rip), %rax
	ret
.L32:
	leaq	.LC16(%rip), %rax
	ret
.L33:
	leaq	.LC17(%rip), %rax
	ret
.L34:
	leaq	.LC18(%rip), %rax
	ret
.L35:
	leaq	.LC19(%rip), %rax
	ret
.L36:
	leaq	.LC20(%rip), %rax
	ret
.L37:
	leaq	.LC21(%rip), %rax
	ret
.L38:
	leaq	.LC22(%rip), %rax
	ret
.L39:
	leaq	.LC23(%rip), %rax
	ret
.L40:
	leaq	.LC24(%rip), %rax
	ret
.L41:
	leaq	.LC25(%rip), %rax
	ret
.L42:
	leaq	.LC26(%rip), %rax
	ret
.L43:
	leaq	.LC27(%rip), %rax
	ret
.L44:
	leaq	.LC28(%rip), %rax
	ret
.L45:
	leaq	.LC29(%rip), %rax
	ret
.L46:
	leaq	.LC30(%rip), %rax
	ret
.L47:
	leaq	.LC31(%rip), %rax
	ret
.L48:
	leaq	.LC32(%rip), %rax
	ret
.L49:
	leaq	.LC33(%rip), %rax
	ret
.L50:
	leaq	.LC34(%rip), %rax
	ret
.L51:
	leaq	.LC35(%rip), %rax
	ret
.L52:
	leaq	.LC36(%rip), %rax
	ret
.L53:
	leaq	.LC37(%rip), %rax
	ret
.L54:
	leaq	.LC38(%rip), %rax
	ret
.L55:
	leaq	.LC39(%rip), %rax
	ret
.L56:
	leaq	.LC40(%rip), %rax
	ret
.L57:
	leaq	.LC41(%rip), %rax
	ret
.L58:
	leaq	.LC42(%rip), %rax
	ret
.L59:
	leaq	.LC43(%rip), %rax
	ret
.L60:
	leaq	.LC44(%rip), %rax
	ret
.L61:
	leaq	.LC45(%rip), %rax
	ret
.L62:
	leaq	.LC46(%rip), %rax
	ret
.L63:
	leaq	.LC47(%rip), %rax
	ret
.L64:
	leaq	.LC48(%rip), %rax
	ret
.L65:
	leaq	.LC49(%rip), %rax
	ret
.L66:
	leaq	.LC50(%rip), %rax
	ret
.L67:
	leaq	.LC51(%rip), %rax
	ret
.L68:
	leaq	.LC52(%rip), %rax
	ret
.L69:
	leaq	.LC53(%rip), %rax
	ret
.L70:
	leaq	.LC54(%rip), %rax
	ret
.L71:
	leaq	.LC55(%rip), %rax
	ret
.L72:
	leaq	.LC56(%rip), %rax
	ret
.L73:
	leaq	.LC57(%rip), %rax
	ret
.L74:
	leaq	.LC58(%rip), %rax
	ret
.L75:
	leaq	.LC59(%rip), %rax
	ret
.L76:
	leaq	.LC60(%rip), %rax
	ret
.L77:
	leaq	.LC61(%rip), %rax
	ret
.L78:
	leaq	.LC62(%rip), %rax
	ret
.L84:
	leaq	.LC65(%rip), %rax
	ret
.L85:
	leaq	.LC66(%rip), %rax
	ret
.L86:
	leaq	.LC67(%rip), %rax
	ret
.L87:
	leaq	.LC68(%rip), %rax
	ret
.L88:
	leaq	.LC69(%rip), %rax
	ret
.L89:
	leaq	.LC70(%rip), %rax
	ret
.L90:
	leaq	.LC71(%rip), %rax
	ret
.L91:
	leaq	.LC72(%rip), %rax
	ret
.L92:
	leaq	.LC73(%rip), %rax
	ret
.L82:
	leaq	.LC77(%rip), %rax
	ret
.L96:
	leaq	.LC81(%rip), %rax
	ret
.L97:
	leaq	.LC82(%rip), %rax
	ret
.L98:
	leaq	.LC83(%rip), %rax
	ret
.L99:
	leaq	.LC84(%rip), %rax
	ret
.L80:
	leaq	.LC85(%rip), %rax
	ret
.L100:
	leaq	.LC86(%rip), %rax
	ret
.L101:
	leaq	.LC87(%rip), %rax
	ret
.L102:
	leaq	.LC88(%rip), %rax
	ret
.L105:
	leaq	.LC90(%rip), %rax
	ret
.L106:
	leaq	.LC63(%rip), %rax
	ret
.L107:
	leaq	.LC75(%rip), %rax
.L10:
	ret
	.cfi_endproc
.LFE2551:
	.size	_Z16clGetErrorStringi, .-_Z16clGetErrorStringi
	.section	.rodata.str1.1
.LC91:
	.string	"Error: %s\n"
.LC92:
	.string	"Error: Failed to create a compute context!\n"
.LC93:
	.string	"Error: Failed to create a command queue!\n"
.LC94:
	.string	"Computing SAXPY on the CPU... Done\n"
.LC95:
	.string	"Build error: %s\n"
.LC96:
	.string	"kernel_upload"
.LC97:
	.string	"Computing SACY on the GPU... Done\n"
.LC99:
	.string	"Error en %d i en dim %d, resultado %f i %f\n"
.LC100:
	.string	"Comparing the output for each implementation... Incorrect!\n"
.LC101:
	.string	"Comparing the output for each implementation... Correct!\n"
.LC102:
	.string	"NUM_PARTICLES: %d\n"
.LC103:
	.string	"BLOCK_SIZE: %d\n"
.LC104:
	.string	"NUM_ITERATIONS: %d\n"
.LC105:
	.string	"Execution time on CPU: %f\n"
.LC106:
	.string	"Execution time on GPU: %f\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2550:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leaq	-479232(%rsp), %r11
	.cfi_def_cfa 11, 479288
.LPSRL0:
	subq	$4096, %rsp
	orq	$0, (%rsp)
	cmpq	%r11, %rsp
	jne	.LPSRL0
	.cfi_def_cfa_register 7
	subq	$2936, %rsp
	.cfi_def_cfa_offset 482224
	xorl	%esi, %esi
	xorl	%edi, %edi
	movq	%fs:40, %rax
	movq	%rax, 482152(%rsp)
	xorl	%eax, %eax
	leaq	60(%rsp), %rdx
	movl	$10000, 52(%rsp)
	call	clGetPlatformIDs@PLT
	movl	%eax, 64(%rsp)
	testl	%eax, %eax
	je	.L114
	movl	%eax, %edi
	movl	$2, %esi
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L114:
	movl	60(%rsp), %edi
	movq	%rdi, %rbx
	salq	$3, %rdi
	call	malloc@PLT
	movl	%ebx, %edi
	xorl	%edx, %edx
	movq	%rax, %rsi
	movq	%rax, %r13
	call	clGetPlatformIDs@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L115
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L115:
	movq	0(%r13), %rdi
	leaq	68(%rsp), %r8
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	movl	$4, %esi
	call	clGetDeviceIDs@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L116
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L116:
	movl	68(%rsp), %edi
	movq	%rdi, %rbx
	salq	$3, %rdi
	call	malloc@PLT
	movq	0(%r13), %rdi
	xorl	%r8d, %r8d
	movl	%ebx, %edx
	movq	%rax, %rcx
	movl	$4, %esi
	movq	%rax, %rbp
	call	clGetDeviceIDs@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L117
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L117:
	movl	68(%rsp), %esi
	leaq	64(%rsp), %r12
	xorl	%edi, %edi
	xorl	%r8d, %r8d
	movq	%r12, %r9
	xorl	%ecx, %ecx
	movq	%rbp, %rdx
	call	clCreateContext@PLT
	movl	64(%rsp), %edi
	movq	%rax, 8(%rsp)
	testl	%edi, %edi
	je	.L118
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L118:
	cmpq	$0, 8(%rsp)
	leaq	.LC92(%rip), %rsi
	je	.L219
	movq	8(%rsp), %rdi
	movq	0(%rbp), %rsi
	movq	%r12, %rcx
	xorl	%edx, %edx
	call	clCreateCommandQueue@PLT
	movl	64(%rsp), %edi
	movq	%rax, %rbx
	testl	%edi, %edi
	je	.L121
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L121:
	testq	%rbx, %rbx
	jne	.L122
	leaq	.LC93(%rip), %rsi
.L219:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %eax
	jmp	.L113
.L122:
	movq	8(%rsp), %rdi
	movq	%r12, %r8
	xorl	%ecx, %ecx
	movl	$1, %esi
	leaq	mykernel(%rip), %rdx
	call	clCreateProgramWithSource@PLT
	movl	64(%rsp), %edi
	movq	%rax, 16(%rsp)
	testl	%edi, %edi
	je	.L124
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L124:
	leaq	104(%rsp), %r14
	movq	%r14, %rdi
	call	_Z10init_ArrayP8particle
	movq	8(%rsp), %rdi
	movq	%r12, %r8
	xorl	%ecx, %ecx
	movl	$240000, %edx
	movl	$1, %esi
	call	clCreateBuffer@PLT
	movl	64(%rsp), %edi
	movq	%rax, 72(%rsp)
	testl	%edi, %edi
	je	.L125
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L125:
	pushq	%rcx
	.cfi_def_cfa_offset 482232
	movq	%rbx, %rdi
	movq	%r14, %r9
	movl	$240000, %r8d
	pushq	$0
	.cfi_def_cfa_offset 482240
	xorl	%ecx, %ecx
	movl	$1, %edx
	pushq	$0
	.cfi_def_cfa_offset 482248
	pushq	$0
	.cfi_def_cfa_offset 482256
	movq	104(%rsp), %rsi
	call	clEnqueueWriteBuffer@PLT
	movl	%eax, 96(%rsp)
	movl	%eax, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 482224
	testl	%eax, %eax
	je	.L126
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L126:
	call	_Z9cpuSecondv
	xorl	%esi, %esi
	movl	$10000, %edi
	movsd	%xmm0, 32(%rsp)
.L129:
	movq	%r14, %rcx
	movl	%esi, %r9d
	movl	%esi, %r10d
	movl	%esi, %r8d
	leal	30000(%rsi), %r11d
.L127:
	cmpl	%r11d, %r8d
	je	.L221
	movl	%r8d, %eax
	addq	$24, %rcx
	addl	$3, %r8d
	cltd
	idivl	%edi
	movl	%r10d, %eax
	addl	$4, %r10d
	cvtsi2ssl	%edx, %xmm2
	cltd
	idivl	%edi
	movss	%xmm2, -12(%rcx)
	movl	%r9d, %eax
	addss	-24(%rcx), %xmm2
	addl	$5, %r9d
	movss	%xmm2, -24(%rcx)
	cvtsi2ssl	%edx, %xmm1
	cltd
	idivl	%edi
	movss	%xmm1, -8(%rcx)
	addss	-20(%rcx), %xmm1
	movss	%xmm1, -20(%rcx)
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, -4(%rcx)
	addss	-16(%rcx), %xmm0
	movss	%xmm0, -16(%rcx)
	jmp	.L127
.L221:
	incl	%esi
	cmpl	$1000, %esi
	jne	.L129
	movl	$10000, 56(%rsp)
	call	_Z9cpuSecondv
	leaq	.LC94(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movsd	%xmm0, 24(%rsp)
	call	__printf_chk@PLT
	movq	16(%rsp), %rdi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	xorl	%ecx, %ecx
	movq	%rbp, %rdx
	movl	$1, %esi
	call	clBuildProgram@PLT
	movl	%eax, 64(%rsp)
	testl	%eax, %eax
	je	.L130
	movq	0(%rbp), %rsi
	movq	16(%rsp), %rdi
	leaq	80(%rsp), %r9
	leaq	480104(%rsp), %rbx
	movq	%rbx, %r8
	movl	$2048, %ecx
	movl	$4483, %edx
	call	clGetProgramBuildInfo@PLT
	movq	%rbx, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	movq	stderr(%rip), %rdi
	leaq	.LC95(%rip), %rdx
	call	__fprintf_chk@PLT
	jmp	.L131
.L130:
	movq	16(%rsp), %rdi
	movq	%r12, %rdx
	leaq	.LC96(%rip), %rsi
	call	clCreateKernel@PLT
	movl	64(%rsp), %edi
	movq	%rax, %r12
	testl	%edi, %edi
	je	.L132
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L132:
	movl	$8, %edx
	movl	$64, %eax
.L133:
	addl	%eax, %eax
	decl	%edx
	jne	.L133
	movq	%r12, %rdi
	leaq	52(%rsp), %rcx
	movl	$4, %edx
	movl	$1, %esi
	movq	%rax, 88(%rsp)
	movq	$64, 96(%rsp)
	call	clSetKernelArg@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L134
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L134:
	movq	%r12, %rdi
	leaq	72(%rsp), %rcx
	movl	$8, %edx
	movl	$2, %esi
	call	clSetKernelArg@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L135
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L135:
	call	_Z9cpuSecondv
	leaq	.LC91(%rip), %r15
	xorl	%eax, %eax
	movsd	%xmm0, 40(%rsp)
.L136:
	movl	%eax, 56(%rsp)
	cmpl	$999, %eax
	jg	.L222
	movq	%r12, %rdi
	leaq	56(%rsp), %rcx
	movl	$4, %edx
	xorl	%esi, %esi
	call	clSetKernelArg@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L137
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	movq	%r15, %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L137:
	pushq	%rdx
	.cfi_def_cfa_offset 482232
	movq	%rbx, %rdi
	xorl	%ecx, %ecx
	movl	$1, %edx
	pushq	$0
	.cfi_def_cfa_offset 482240
	movq	%r12, %rsi
	pushq	$0
	.cfi_def_cfa_offset 482248
	pushq	$0
	.cfi_def_cfa_offset 482256
	leaq	128(%rsp), %r9
	leaq	120(%rsp), %r8
	call	clEnqueueNDRangeKernel@PLT
	movl	%eax, 96(%rsp)
	movl	%eax, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 482224
	testl	%eax, %eax
	je	.L138
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	movq	%r15, %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L138:
	movq	%rbx, %rdi
	call	clFlush@PLT
	movq	%rbx, %rdi
	movl	%eax, 64(%rsp)
	call	clFinish@PLT
	movl	%eax, 64(%rsp)
	movl	56(%rsp), %eax
	incl	%eax
	jmp	.L136
.L222:
	movq	%rbx, %rdi
	movl	$240000, %r8d
	xorl	%ecx, %ecx
	movl	$1, %edx
	leaq	240104(%rsp), %r15
	pushq	%rax
	.cfi_def_cfa_offset 482232
	pushq	$0
	.cfi_def_cfa_offset 482240
	movq	%r15, %r9
	pushq	$0
	.cfi_def_cfa_offset 482248
	pushq	$0
	.cfi_def_cfa_offset 482256
	movq	104(%rsp), %rsi
	call	clEnqueueReadBuffer@PLT
	movl	%eax, 96(%rsp)
	movl	%eax, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 482224
	testl	%eax, %eax
	je	.L140
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L140:
	call	_Z9cpuSecondv
	subsd	40(%rsp), %xmm0
	leaq	.LC97(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	movsd	%xmm0, 40(%rsp)
	call	__printf_chk@PLT
	movsd	.LC0(%rip), %xmm2
	xorl	%eax, %eax
	movss	.LC98(%rip), %xmm1
.L141:
	imulq	$24, %rax, %rdx
	xorl	%ecx, %ecx
.L145:
	movss	(%r15,%rdx), %xmm0
	subss	(%r14,%rdx), %xmm0
	andps	%xmm1, %xmm0
	cvtss2sd	%xmm0, %xmm0
	comisd	%xmm2, %xmm0
	jbe	.L217
	movl	%eax, %edx
	cltq
	leaq	.LC99(%rip), %rsi
	movl	$2, %edi
	imulq	$24, %rax, %rax
	cvtss2sd	240108(%rsp,%rax), %xmm0
	cvtss2sd	108(%rsp,%rax), %xmm1
	movb	$2, %al
	call	__printf_chk@PLT
	leaq	.LC100(%rip), %rsi
	jmp	.L220
.L217:
	incl	%ecx
	addq	$4, %rdx
	cmpl	$3, %ecx
	jne	.L145
	incq	%rax
	cmpq	$10000, %rax
	jne	.L141
	leaq	.LC101(%rip), %rsi
.L220:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$10000, %edx
	leaq	.LC102(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$64, %edx
	leaq	.LC103(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1000, %edx
	leaq	.LC104(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movsd	24(%rsp), %xmm0
	subsd	32(%rsp), %xmm0
	movb	$1, %al
	leaq	.LC105(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movsd	40(%rsp), %xmm0
	leaq	.LC106(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	72(%rsp), %rdi
	call	clReleaseMemObject@PLT
	movl	64(%rsp), %edi
	testl	%edi, %edi
	je	.L147
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L147:
	movq	16(%rsp), %rdi
	call	clReleaseProgram@PLT
	movl	64(%rsp), %edi
	testl	%edi, %edi
	je	.L148
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L148:
	movq	%r12, %rdi
	call	clReleaseKernel@PLT
	movl	64(%rsp), %edi
	testl	%edi, %edi
	je	.L149
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L149:
	movq	%rbx, %rdi
	call	clReleaseCommandQueue@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L150
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L150:
	movq	8(%rsp), %rdi
	call	clReleaseContext@PLT
	movl	%eax, 64(%rsp)
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L151
	call	_Z16clGetErrorStringi
	movq	stderr(%rip), %rdi
	leaq	.LC91(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L151:
	movq	%r13, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
.L131:
	xorl	%eax, %eax
.L113:
	movq	482152(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L152
	call	__stack_chk_fail@PLT
.L152:
	addq	$482168, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2550:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2577:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2577:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	mykernel
	.section	.rodata.str1.1
.LC107:
	.ascii	"typedef struct{                                         \nfl"
	.ascii	"oat position[3];                                    \nfloat "
	.ascii	"velocity[3];                                   \n}particle; "
	.ascii	"                                                      \n__ke"
	.ascii	"rnel                                                 \nvoid "
	.ascii	"kernel_upload(                                       \n   co"
	.ascii	"nst int iter,                                          \n   "
	.ascii	"const int nump,                                        \n   "
	.ascii	"__global particle* X)                                    \n{"
	.ascii	"int index = get_global_id(0);                           \nif"
	.ascii	"(index < nump) {                                            "
	.ascii	"    \n   X[index].velocity[0] = (3*index + iter) % nump;    "
	.ascii	"         \n   X[index].velocity[1] = (4*index + iter) % nump"
	.ascii	";            \n   X[index].velocity[2] = (5*index + iter) % "
	.ascii	"nump;             \n   X[index].position[0] = X[index].posit"
	.ascii	"ion[0] + X[index].velocity[0];             \n   X[index].pos"
	.ascii	"ition[1] = X[in"
	.string	"dex].position[1] + X[index].velocity[1];             \n   X[index].position[2] = X[index].position[2] + X[index].velocity[2];             \n}                                                          \n}                                                        \n"
	.section	.data.rel.local,"aw"
	.align 8
	.type	mykernel, @object
	.size	mykernel, 8
mykernel:
	.quad	.LC107
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-1598689907
	.long	1051772663
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC98:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.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:

// Template file for the OpenCL Assignment 4


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>
#include <math.h>
#include <sys/time.h>


#define NUM_PARTICLES 10000
#define NUM_ITERATIONS 1000
#define BLOCK_SIZE 64
#define error 1e-6
// This is a macro for checking the error variable:>.
#define CHK_ERROR(err) if (err != CL_SUCCESS) fprintf(stderr,"Error: %s\n",clGetErrorString(err));

// A errorCode to string converter (forward declaration)
const char* clGetErrorString(int);




const char *mykernel = 
"typedef struct{                                         \n"
    "float position[3];                                    \n"
    "float velocity[3];                                   \n"
"}particle;                                                       \n"
"__kernel                                                 \n" 
"void kernel_upload(                                       \n"
"   const int iter,                                          \n"
"   const int nump,                                        \n"    
"   __global particle* X)                                    \n"                                      
"{int index = get_global_id(0);                           \n"
"if(index < nump) {                                                \n"
"   X[index].velocity[0] = (3*index + iter) % nump;             \n"
"   X[index].velocity[1] = (4*index + iter) % nump;            \n"
"   X[index].velocity[2] = (5*index + iter) % nump;             \n"
"   X[index].position[0] = X[index].position[0] + X[index].velocity[0];             \n"    
"   X[index].position[1] = X[index].position[1] + X[index].velocity[1];             \n"
"   X[index].position[2] = X[index].position[2] + X[index].velocity[2];             \n"
"}                                                          \n"
"}                                                        \n";

typedef struct{
    float position[3];
    float velocity[3];
}particle;

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp,NULL);
    return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
}

void init_Array(particle *particulas){
    for(int i = 0; i < NUM_PARTICLES; i++){
        particulas[i].position[0] = rand() % 1000;
        particulas[i].position[1] = rand() % 1000;
        particulas[i].position[2] = rand() % 1000;

        particulas[i].velocity[0] = rand() % 1000;
        particulas[i].velocity[1] = rand() % 1000;
        particulas[i].velocity[2] = rand() % 1000;
    }
}


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

    //All the variables used
    particle X[NUM_PARTICLES],  results[NUM_PARTICLES];
    float array = NUM_PARTICLES * sizeof(particle);
    int count = NUM_PARTICLES;
    bool correct = true;
    int i, j;

    cl_platform_id * platforms; cl_uint     n_platform;

    // Find OpenCL Platforms
    cl_int err = clGetPlatformIDs(0, NULL, &n_platform); CHK_ERROR(err);
    platforms = (cl_platform_id *) malloc(sizeof(cl_platform_id)*n_platform);
    err = clGetPlatformIDs(n_platform, platforms, NULL); CHK_ERROR(err);

    // Find and sort devices
    cl_device_id *device_list; cl_uint n_devices;
    err = clGetDeviceIDs( platforms[0], CL_DEVICE_TYPE_GPU, 0,NULL, &n_devices);CHK_ERROR(err);
    device_list = (cl_device_id *) malloc(sizeof(cl_device_id)*n_devices);
    err = clGetDeviceIDs( platforms[0],CL_DEVICE_TYPE_GPU, n_devices, device_list, NULL);CHK_ERROR(err);
    
    // Create and initialize an OpenCL context
    cl_context context = clCreateContext( NULL, n_devices, device_list, NULL, NULL, &err);CHK_ERROR(err);

    if (!context) {
        printf("Error: Failed to create a compute context!\n");
        return EXIT_FAILURE;
    }

    // Create a command queue
    cl_command_queue cmd_queue = clCreateCommandQueue(context, device_list[0], 0, &err);CHK_ERROR(err);

    if (!cmd_queue) {
        printf("Error: Failed to create a command queue!\n");
        return EXIT_FAILURE;
    }


    /* Insert your own code here */

    //program
    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&mykernel, NULL, &err);CHK_ERROR(err);


    //Initialize values of the arrays
 
    init_Array(X);

    //Array on GPU part
    cl_mem dx = clCreateBuffer(context, CL_MEM_READ_WRITE, array, NULL, &err);CHK_ERROR(err);


    //Copy the values of X in dx
    err = clEnqueueWriteBuffer(cmd_queue, dx, CL_TRUE, 0, array, X, 0, NULL, NULL);CHK_ERROR(err);


    //Cpu part
    double start = cpuSecond();
    for(j = 0; j < NUM_ITERATIONS; j++) {
        for(i = 0; i < NUM_PARTICLES; i++){
            X[i].velocity[0] = (3*i + j) % NUM_PARTICLES;             
            X[i].velocity[1] = (4*i + j) % NUM_PARTICLES;            
            X[i].velocity[2] = (5*i + j) % NUM_PARTICLES;             
            X[i].position[0] = X[i].position[0] + X[i].velocity[0];                 
            X[i].position[1] = X[i].position[1] + X[i].velocity[1];            
            X[i].position[2] = X[i].position[2] + X[i].velocity[2]; 
        }
    }
    double totalCpu = cpuSecond() - start;        

    printf("Computing SAXPY on the CPU... Done\n");
    
    //build the program 
    err = clBuildProgram(program, 1, device_list, NULL, NULL, NULL);
    if (err != CL_SUCCESS){
        size_t len;
        char buffer[2048];
        clGetProgramBuildInfo(program, device_list[0], CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        fprintf(stderr, "Build error: %s\n", buffer);
        return 0;
    }

    //Create the kernel function
    cl_kernel kernel = clCreateKernel(program, "kernel_upload", &err);CHK_ERROR(err);
    


    //More argument, number of workitems we are going to use
    unsigned int mult = BLOCK_SIZE;
    unsigned int size = BLOCK_SIZE;
    while(mult < NUM_PARTICLES){

        mult = mult * 2;
    }
    size_t n_workitem[1] = {mult};
    size_t workgroup_size[1] = {size};

    err = clSetKernelArg(kernel, 1, sizeof(int), (void*) &count);CHK_ERROR(err);
    err = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &dx);CHK_ERROR(err);

    
    double st = cpuSecond();
    for(i = 0; i < NUM_ITERATIONS; i++){

        //Arguments for the Kernel function
        err = clSetKernelArg(kernel, 0, sizeof(int), (void*) &i);CHK_ERROR(err);



        //Launch kernel
        err = clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL, n_workitem, workgroup_size, 0, NULL, NULL);CHK_ERROR(err);


        //Wait for everything to finish
        err = clFlush(cmd_queue);
        err = clFinish(cmd_queue);
    }
    err = clEnqueueReadBuffer(cmd_queue, dx, CL_TRUE, 0, array, results, 0, NULL, NULL);CHK_ERROR(err);
    double timeGpu = cpuSecond() - st;
    printf("Computing SACY on the GPU... Done\n");

    for(int i = 0; i < NUM_PARTICLES; i++){
        for(int dim = 0; dim < 3; dim++){
            if(abs(results[i].position[dim] - X[i].position[dim]) > error ){
                correct = false;
                printf("Error en %d i en dim %d, resultado %f i %f\n",i,dim,results[i].position[1],X[i].position[1]);
                break;
            }
        }
        if(!correct) {break;}
    }

    if (correct) {
        printf("Comparing the output for each implementation... Correct!\n");
    }
    else {
        printf("Comparing the output for each implementation... Incorrect!\n");
    }

    printf("NUM_PARTICLES: %d\n", NUM_PARTICLES);
    printf("BLOCK_SIZE: %d\n", BLOCK_SIZE);
    printf("NUM_ITERATIONS: %d\n", NUM_ITERATIONS);
    printf("Execution time on CPU: %f\n", totalCpu);
    printf("Execution time on GPU: %f\n", timeGpu);


    // Finally, release all that we have allocated.
    clReleaseMemObject(dx);CHK_ERROR(err);
    clReleaseProgram(program);CHK_ERROR(err);
    clReleaseKernel(kernel);CHK_ERROR(err);
    err = clReleaseCommandQueue(cmd_queue);CHK_ERROR(err);
    err = clReleaseContext(context);CHK_ERROR(err);
    free(platforms);
    free(device_list);
    
    return 0;
}



// The source for this particular version is from: https://stackoverflow.com/questions/24326432/convenient-way-to-show-opencl-error-codes
const char* clGetErrorString(int errorCode) {
    switch (errorCode) {
    case 0: return "CL_SUCCESS";
    case -1: return "CL_DEVICE_NOT_FOUND";
    case -2: return "CL_DEVICE_NOT_AVAILABLE";
    case -3: return "CL_COMPILER_NOT_AVAILABLE";
    case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
    case -5: return "CL_OUT_OF_RESOURCES";
    case -6: return "CL_OUT_OF_HOST_MEMORY";
    case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE";
    case -8: return "CL_MEM_COPY_OVERLAP";
    case -9: return "CL_IMAGE_FORMAT_MISMATCH";
    case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
    case -12: return "CL_MAP_FAILURE";
    case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
    case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
    case -15: return "CL_COMPILE_PROGRAM_FAILURE";
    case -16: return "CL_LINKER_NOT_AVAILABLE";
    case -17: return "CL_LINK_PROGRAM_FAILURE";
    case -18: return "CL_DEVICE_PARTITION_FAILED";
    case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
    case -30: return "CL_INVALID_VALUE";
    case -31: return "CL_INVALID_DEVICE_TYPE";
    case -32: return "CL_INVALID_PLATFORM";
    case -33: return "CL_INVALID_DEVICE";
    case -34: return "CL_INVALID_CONTEXT";
    case -35: return "CL_INVALID_QUEUE_PROPERTIES";
    case -36: return "CL_INVALID_COMMAND_QUEUE";
    case -37: return "CL_INVALID_HOST_PTR";
    case -38: return "CL_INVALID_MEM_OBJECT";
    case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
    case -40: return "CL_INVALID_IMAGE_SIZE";
    case -41: return "CL_INVALID_SAMPLER";
    case -42: return "CL_INVALID_BINARY";
    case -43: return "CL_INVALID_BUILD_OPTIONS";
    case -44: return "CL_INVALID_PROGRAM";
    case -45: return "CL_INVALID_PROGRAM_EXECUTABLE";
    case -46: return "CL_INVALID_KERNEL_NAME";
    case -47: return "CL_INVALID_KERNEL_DEFINITION";
    case -48: return "CL_INVALID_KERNEL";
    case -49: return "CL_INVALID_ARG_INDEX";
    case -50: return "CL_INVALID_ARG_VALUE";
    case -51: return "CL_INVALID_ARG_SIZE";
    case -52: return "CL_INVALID_KERNEL_ARGS";
    case -53: return "CL_INVALID_WORK_DIMENSION";
    case -54: return "CL_INVALID_WORK_GROUP_SIZE";
    case -55: return "CL_INVALID_WORK_ITEM_SIZE";
    case -56: return "CL_INVALID_GLOBAL_OFFSET";
    case -57: return "CL_INVALID_EVENT_WAIT_LIST";
    case -58: return "CL_INVALID_EVENT";
    case -59: return "CL_INVALID_OPERATION";
    case -60: return "CL_INVALID_GL_OBJECT";
    case -61: return "CL_INVALID_BUFFER_SIZE";
    case -62: return "CL_INVALID_MIP_LEVEL";
    case -63: return "CL_INVALID_GLOBAL_WORK_SIZE";
    case -64: return "CL_INVALID_PROPERTY";
    case -65: return "CL_INVALID_IMAGE_DESCRIPTOR";
    case -66: return "CL_INVALID_COMPILER_OPTIONS";
    case -67: return "CL_INVALID_LINKER_OPTIONS";
    case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT";
    case -69: return "CL_INVALID_PIPE_SIZE";
    case -70: return "CL_INVALID_DEVICE_QUEUE";
    case -71: return "CL_INVALID_SPEC_ID";
    case -72: return "CL_MAX_SIZE_RESTRICTION_EXCEEDED";
    case -1002: return "CL_INVALID_D3D10_DEVICE_KHR";
    case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR";
    case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
    case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
    case -1006: return "CL_INVALID_D3D11_DEVICE_KHR";
    case -1007: return "CL_INVALID_D3D11_RESOURCE_KHR";
    case -1008: return "CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR";
    case -1009: return "CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR";
    case -1010: return "CL_INVALID_DX9_MEDIA_ADAPTER_KHR";
    case -1011: return "CL_INVALID_DX9_MEDIA_SURFACE_KHR";
    case -1012: return "CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR";
    case -1013: return "CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR";
    case -1093: return "CL_INVALID_EGL_OBJECT_KHR";
    case -1092: return "CL_EGL_RESOURCE_NOT_ACQUIRED_KHR";
    case -1001: return "CL_PLATFORM_NOT_FOUND_KHR";
    case -1057: return "CL_DEVICE_PARTITION_FAILED_EXT";
    case -1058: return "CL_INVALID_PARTITION_COUNT_EXT";
    case -1059: return "CL_INVALID_PARTITION_NAME_EXT";
    case -1094: return "CL_INVALID_ACCELERATOR_INTEL";
    case -1095: return "CL_INVALID_ACCELERATOR_TYPE_INTEL";
    case -1096: return "CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL";
    case -1097: return "CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL";
    case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
    case -1098: return "CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL";
    case -1099: return "CL_INVALID_VA_API_MEDIA_SURFACE_INTEL";
    case -1100: return "CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL";
    case -1101: return "CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL";
    default: return "CL_UNKNOWN_ERROR";
    }
}


	.text
	.file	"exercise_3.hip"
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z9cpuSecondv
.LCPI0_0:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
	.text
	.globl	_Z9cpuSecondv
	.type	_Z9cpuSecondv,@function
_Z9cpuSecondv:                          # @_Z9cpuSecondv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	movq	%rsp, %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	cvtsi2sdq	(%rbx), %xmm1
	cvtsi2sdq	8(%rbx), %xmm0
	mulsd	.LCPI0_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z9cpuSecondv, .Lfunc_end0-_Z9cpuSecondv
	.cfi_endproc
                                        # -- End function
	.globl	_Z10init_ArrayP8particle        # -- Begin function _Z10init_ArrayP8particle
	.type	_Z10init_ArrayP8particle,@function
_Z10init_ArrayP8particle:               # @_Z10init_ArrayP8particle
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	movl	$20, %r14d
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, -20(%rbx,%r14)
	callq	rand
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, -16(%rbx,%r14)
	callq	rand
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, -12(%rbx,%r14)
	callq	rand
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, -8(%rbx,%r14)
	callq	rand
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, -4(%rbx,%r14)
	callq	rand
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, (%rbx,%r14)
	addq	$24, %r14
	cmpq	$240020, %r14                   # imm = 0x3A994
	jne	.LBB1_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z10init_ArrayP8particle, .Lfunc_end1-_Z10init_ArrayP8particle
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI2_0:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI2_1:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
.LCPI2_2:
	.zero	16
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$480168, %rsp                   # imm = 0x753A8
	.cfi_def_cfa_offset 480224
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$10000, 84(%rsp)                # imm = 0x2710
	leaq	80(%rsp), %rdx
	xorl	%edi, %edi
	xorl	%esi, %esi
	callq	clGetPlatformIDs
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_1
.LBB2_2:
	movl	80(%rsp), %ebx
	movq	%rbx, %rdi
	shlq	$3, %rdi
	callq	malloc
	movq	%rax, %r13
	movl	%ebx, %edi
	movq	%rax, %rsi
	xorl	%edx, %edx
	callq	clGetPlatformIDs
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_3
.LBB2_4:
	movq	(%r13), %rdi
	leaq	60(%rsp), %r8
	movl	$4, %esi
	xorl	%edx, %edx
	xorl	%ecx, %ecx
	callq	clGetDeviceIDs
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_5
.LBB2_6:
	movl	60(%rsp), %ebx
	movq	%rbx, %rdi
	shlq	$3, %rdi
	callq	malloc
	movq	%rax, %r15
	movq	(%r13), %rdi
	movl	$4, %esi
	movl	%ebx, %edx
	movq	%rax, %rcx
	xorl	%r8d, %r8d
	callq	clGetDeviceIDs
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_7
.LBB2_8:
	movl	60(%rsp), %esi
	leaq	36(%rsp), %rbx
	xorl	%edi, %edi
	movq	%r15, %rdx
	xorl	%ecx, %ecx
	xorl	%r8d, %r8d
	movq	%rbx, %r9
	callq	clCreateContext
	movq	%rax, %rbp
	movl	(%rbx), %edi
	testl	%edi, %edi
	jne	.LBB2_9
.LBB2_10:
	testq	%rbp, %rbp
	je	.LBB2_11
# %bb.13:
	movq	(%r15), %rsi
	leaq	36(%rsp), %rbx
	movq	%rbp, %rdi
	xorl	%edx, %edx
	movq	%rbx, %rcx
	callq	clCreateCommandQueue
	movq	%rax, %r12
	movl	(%rbx), %edi
	testl	%edi, %edi
	jne	.LBB2_14
.LBB2_15:
	testq	%r12, %r12
	je	.LBB2_16
# %bb.17:
	leaq	36(%rsp), %rbx
	movl	$mykernel, %edx
	movq	%rbp, %rdi
	movl	$1, %esi
	xorl	%ecx, %ecx
	movq	%rbx, %r8
	callq	clCreateProgramWithSource
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	movl	(%rbx), %edi
	testl	%edi, %edi
	jne	.LBB2_18
.LBB2_19:
	leaq	160(%rsp), %rdi
	callq	_Z10init_ArrayP8particle
	leaq	36(%rsp), %r14
	movl	$1, %esi
	movl	$240000, %edx                   # imm = 0x3A980
	movq	%rbp, 120(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	xorl	%ecx, %ecx
	movq	%r14, %r8
	callq	clCreateBuffer
	movq	%rax, %rbx
	movq	%rax, 72(%rsp)
	movl	(%r14), %edi
	testl	%edi, %edi
	jne	.LBB2_20
.LBB2_21:
	movq	%r15, 96(%rsp)                  # 8-byte Spill
	xorps	%xmm0, %xmm0
	movups	%xmm0, 8(%rsp)
	movl	$0, (%rsp)
	leaq	160(%rsp), %r9
	movl	$240000, %r8d                   # imm = 0x3A980
	movq	%r12, %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	xorl	%ecx, %ecx
	callq	clEnqueueWriteBuffer
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_22
.LBB2_23:
	xorl	%r15d, %r15d
	leaq	240160(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	cvtsi2sdq	8(%rbx), %xmm3
	movq	(%rbx), %r14
	mulsd	.LCPI2_0(%rip), %xmm3
	movl	$3518437209, %eax               # imm = 0xD1B71759
.LBB2_24:                               # %.preheader126
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_25 Depth 2
	xorl	%ecx, %ecx
	movl	%r15d, %edx
	movl	%r15d, %esi
	movl	%r15d, %edi
.LBB2_25:                               #   Parent Loop BB2_24 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%edi, %r8d
	imulq	%rax, %r8
	shrq	$45, %r8
	imull	$10000, %r8d, %r8d              # imm = 0x2710
	movl	%edi, %r9d
	subl	%r8d, %r9d
	movl	%esi, %r8d
	imulq	%rax, %r8
	shrq	$45, %r8
	imull	$10000, %r8d, %r8d              # imm = 0x2710
	movl	%esi, %r10d
	subl	%r8d, %r10d
	movl	%edx, %r8d
	imulq	%rax, %r8
	shrq	$45, %r8
	imull	$10000, %r8d, %r8d              # imm = 0x2710
	movl	%edx, %r11d
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r10d, %xmm0
	subl	%r8d, %r11d
	xorps	%xmm1, %xmm1
	cvtsi2ss	%r9d, %xmm1
	movss	%xmm0, 172(%rsp,%rcx,8)
	xorps	%xmm2, %xmm2
	cvtsi2ss	%r11d, %xmm2
	movss	%xmm1, 176(%rsp,%rcx,8)
	movss	%xmm2, 180(%rsp,%rcx,8)
	addss	160(%rsp,%rcx,8), %xmm0
	movss	%xmm0, 160(%rsp,%rcx,8)
	addss	164(%rsp,%rcx,8), %xmm1
	movss	%xmm1, 164(%rsp,%rcx,8)
	addss	168(%rsp,%rcx,8), %xmm2
	movss	%xmm2, 168(%rsp,%rcx,8)
	addl	$4, %edi
	addl	$3, %esi
	addl	$5, %edx
	addq	$3, %rcx
	cmpq	$30000, %rcx                    # imm = 0x7530
	jne	.LBB2_25
# %bb.26:                               #   in Loop: Header=BB2_24 Depth=1
	incq	%r15
	cmpq	$1000, %r15                     # imm = 0x3E8
	jne	.LBB2_24
# %bb.27:
	movsd	%xmm3, 40(%rsp)                 # 8-byte Spill
	movl	$10000, 56(%rsp)                # imm = 0x2710
	leaq	240160(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	xorps	%xmm0, %xmm0
	cvtsi2sdq	(%rbx), %xmm0
	movsd	%xmm0, 64(%rsp)                 # 8-byte Spill
	xorps	%xmm0, %xmm0
	cvtsi2sdq	8(%rbx), %xmm0
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movq	88(%rsp), %r15                  # 8-byte Reload
	movq	%r15, %rdi
	movl	$1, %esi
	movq	96(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rdx
	xorl	%ecx, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	clBuildProgram
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_28
# %bb.29:
	movq	%r13, 112(%rsp)                 # 8-byte Spill
	leaq	36(%rsp), %rbx
	movl	$.L.str.6, %esi
	movq	%r15, %rdi
	movq	%rbx, %rdx
	callq	clCreateKernel
	movq	%rax, %r13
	movl	(%rbx), %edi
	testl	%edi, %edi
	jne	.LBB2_30
.LBB2_31:
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r14, %xmm0
	movsd	%xmm0, 104(%rsp)                # 8-byte Spill
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	mulsd	.LCPI2_0(%rip), %xmm0
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	movq	$16384, 136(%rsp)               # imm = 0x4000
	movq	$64, 128(%rsp)
	leaq	84(%rsp), %rcx
	movl	$4, %edx
	movq	%r13, %rdi
	movl	$1, %esi
	callq	clSetKernelArg
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_32
.LBB2_33:
	movsd	40(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	addsd	104(%rsp), %xmm0                # 8-byte Folded Reload
	movsd	%xmm0, 40(%rsp)                 # 8-byte Spill
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	addsd	64(%rsp), %xmm0                 # 8-byte Folded Reload
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	leaq	72(%rsp), %rcx
	movl	$8, %edx
	movq	%r13, %rdi
	movl	$2, %esi
	callq	clSetKernelArg
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_34
.LBB2_35:
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	subsd	40(%rsp), %xmm0                 # 8-byte Folded Reload
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	leaq	240160(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	xorps	%xmm0, %xmm0
	cvtsi2sdq	8(%rbx), %xmm0
	movq	(%rbx), %rax
	movq	%rax, 64(%rsp)                  # 8-byte Spill
	mulsd	.LCPI2_0(%rip), %xmm0
	movsd	%xmm0, 40(%rsp)                 # 8-byte Spill
	leaq	56(%rsp), %rbx
	movl	$0, (%rbx)
	leaq	136(%rsp), %r15
	leaq	128(%rsp), %r14
.LBB2_36:                               # =>This Inner Loop Header: Depth=1
	movl	$4, %edx
	movq	%r13, %rdi
	xorl	%esi, %esi
	movq	%rbx, %rcx
	callq	clSetKernelArg
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_37
.LBB2_38:                               #   in Loop: Header=BB2_36 Depth=1
	xorpd	%xmm0, %xmm0
	movupd	%xmm0, 8(%rsp)
	movl	$0, (%rsp)
	movq	%r12, %rdi
	movq	%r13, %rsi
	movl	$1, %edx
	xorl	%ecx, %ecx
	movq	%r15, %r8
	movq	%r14, %r9
	callq	clEnqueueNDRangeKernel
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_39
.LBB2_40:                               #   in Loop: Header=BB2_36 Depth=1
	movq	%r12, %rdi
	callq	clFlush
	movl	%eax, 36(%rsp)
	movq	%r12, %rdi
	callq	clFinish
	movl	%eax, 36(%rsp)
	movl	56(%rsp), %eax
	leal	1(%rax), %ecx
	movl	%ecx, 56(%rsp)
	cmpl	$999, %eax                      # imm = 0x3E7
	jl	.LBB2_36
	jmp	.LBB2_41
.LBB2_37:                               #   in Loop: Header=BB2_36 Depth=1
	movq	stderr(%rip), %rbp
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbp, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_38
.LBB2_39:                               #   in Loop: Header=BB2_36 Depth=1
	movq	stderr(%rip), %rbp
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbp, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_40
.LBB2_11:
	movl	$.Lstr, %edi
	jmp	.LBB2_12
.LBB2_41:
	xorps	%xmm0, %xmm0
	cvtsi2sdq	64(%rsp), %xmm0         # 8-byte Folded Reload
	movsd	40(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	addsd	%xmm0, %xmm1
	movsd	%xmm1, 40(%rsp)                 # 8-byte Spill
	movq	72(%rsp), %rsi
	xorpd	%xmm0, %xmm0
	movupd	%xmm0, 8(%rsp)
	movl	$0, (%rsp)
	leaq	240160(%rsp), %rbx
	movl	$240000, %r8d                   # imm = 0x3A980
	movq	%r12, %rdi
	movl	$1, %edx
	xorl	%ecx, %ecx
	movq	%rbx, %r9
	callq	clEnqueueReadBuffer
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_42
.LBB2_43:
	xorl	%r14d, %r14d
	leaq	144(%rsp), %r15
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	xorps	%xmm0, %xmm0
	cvtsi2sdq	(%r15), %xmm0
	xorps	%xmm1, %xmm1
	cvtsi2sdq	8(%r15), %xmm1
	mulsd	.LCPI2_0(%rip), %xmm1
	addsd	%xmm0, %xmm1
	subsd	40(%rsp), %xmm1                 # 8-byte Folded Reload
	movsd	%xmm1, 40(%rsp)                 # 8-byte Spill
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movaps	.LCPI2_1(%rip), %xmm0           # xmm0 = [NaN,NaN,NaN,NaN]
	movsd	.LCPI2_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movq	96(%rsp), %r15                  # 8-byte Reload
	leaq	160(%rsp), %rsi
.LBB2_44:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_45 Depth 2
	leaq	(,%r14,8), %rax
	leaq	(%rax,%rax,2), %rcx
	leaq	(%rsp,%rcx), %rax
	addq	$240160, %rax                   # imm = 0x3AA20
	addq	%rsp, %rcx
	addq	$160, %rcx
	xorl	%edx, %edx
.LBB2_45:                               #   Parent Loop BB2_44 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%rbx,%rdx,4), %xmm2            # xmm2 = mem[0],zero,zero,zero
	subss	(%rsi,%rdx,4), %xmm2
	andps	%xmm0, %xmm2
	cvtss2sd	%xmm2, %xmm2
	ucomisd	%xmm1, %xmm2
	ja	.LBB2_46
# %bb.47:                               #   in Loop: Header=BB2_45 Depth=2
	incq	%rdx
	cmpq	$3, %rdx
	jne	.LBB2_45
# %bb.48:                               #   in Loop: Header=BB2_44 Depth=1
	incq	%r14
	addq	$24, %rsi
	addq	$24, %rbx
	cmpq	$10000, %r14                    # imm = 0x2710
	jne	.LBB2_44
# %bb.49:
	movl	$.Lstr.5, %edi
	jmp	.LBB2_50
.LBB2_46:                               # %.thread123
	xorps	%xmm0, %xmm0
	cvtss2sd	4(%rax), %xmm0
	xorps	%xmm1, %xmm1
	cvtss2sd	4(%rcx), %xmm1
	movl	$.L.str.8, %edi
	movl	%r14d, %esi
                                        # kill: def $edx killed $edx killed $rdx
	movb	$2, %al
	callq	printf
	movl	$.Lstr.4, %edi
.LBB2_50:                               # %.loopexit
	callq	puts@PLT
	movl	$.L.str.11, %edi
	movl	$10000, %esi                    # imm = 0x2710
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.12, %edi
	movl	$64, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.13, %edi
	movl	$1000, %esi                     # imm = 0x3E8
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.14, %edi
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	movl	$.L.str.15, %edi
	movsd	40(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	movq	72(%rsp), %rdi
	callq	clReleaseMemObject
	movl	36(%rsp), %edi
	testl	%edi, %edi
	jne	.LBB2_51
.LBB2_52:
	movq	88(%rsp), %rdi                  # 8-byte Reload
	callq	clReleaseProgram
	movl	36(%rsp), %edi
	testl	%edi, %edi
	jne	.LBB2_53
.LBB2_54:
	movq	%r13, %rdi
	callq	clReleaseKernel
	movl	36(%rsp), %edi
	testl	%edi, %edi
	jne	.LBB2_55
.LBB2_56:
	movq	%r12, %rdi
	callq	clReleaseCommandQueue
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	movq	112(%rsp), %r14                 # 8-byte Reload
	jne	.LBB2_57
.LBB2_58:
	movq	120(%rsp), %rdi                 # 8-byte Reload
	callq	clReleaseContext
	movl	%eax, 36(%rsp)
	testl	%eax, %eax
	jne	.LBB2_59
.LBB2_60:
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
.LBB2_61:
	xorl	%eax, %eax
	jmp	.LBB2_62
.LBB2_16:
	movl	$.Lstr.1, %edi
.LBB2_12:
	callq	puts@PLT
	movl	$1, %eax
.LBB2_62:
	addq	$480168, %rsp                   # imm = 0x753A8
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB2_1:
	.cfi_def_cfa_offset 480224
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_2
.LBB2_3:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_4
.LBB2_5:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_6
.LBB2_7:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_8
.LBB2_9:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_10
.LBB2_14:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_15
.LBB2_18:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_19
.LBB2_20:
	movq	stderr(%rip), %r14
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%r14, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_21
.LBB2_22:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_23
.LBB2_28:
	movq	(%rbx), %rsi
	leaq	240160(%rsp), %rbx
	leaq	144(%rsp), %r9
	movl	$2048, %ecx                     # imm = 0x800
	movq	%r15, %rdi
	movl	$4483, %edx                     # imm = 0x1183
	movq	%rbx, %r8
	callq	clGetProgramBuildInfo
	movq	stderr(%rip), %rdi
	movl	$.L.str.5, %esi
	movq	%rbx, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_61
.LBB2_30:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_31
.LBB2_32:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_33
.LBB2_34:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_35
.LBB2_42:
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%r14, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_43
.LBB2_51:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_52
.LBB2_53:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_54
.LBB2_55:
	movq	stderr(%rip), %rbx
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_56
.LBB2_57:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_58
.LBB2_59:
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	_Z16clGetErrorStringi
	movl	$.L.str.1, %esi
	movq	%rbx, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB2_60
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z16clGetErrorStringi           # -- Begin function _Z16clGetErrorStringi
	.type	_Z16clGetErrorStringi,@function
_Z16clGetErrorStringi:                  # @_Z16clGetErrorStringi
	.cfi_startproc
# %bb.0:
                                        # kill: def $edi killed $edi def $rdi
	cmpl	$-73, %edi
	jle	.LBB3_1
# %bb.10:
	addl	$72, %edi
	cmpl	$72, %edi
	ja	.LBB3_99
# %bb.11:
	jmpq	*.LJTI3_2(,%rdi,8)
.LBB3_73:
	movl	$.L.str.77, %eax
	retq
.LBB3_1:
	cmpl	$-1058, %edi                    # imm = 0xFBDE
	jle	.LBB3_2
# %bb.6:
	leal	1013(%rdi), %eax
	cmpl	$13, %eax
	ja	.LBB3_7
# %bb.9:
	jmpq	*.LJTI3_1(,%rax,8)
.LBB3_85:
	movl	$.L.str.89, %eax
	retq
.LBB3_2:
	leal	1101(%rdi), %eax
	cmpl	$9, %eax
	ja	.LBB3_3
# %bb.101:
	jmpq	*.LJTI3_0(,%rax,8)
.LBB3_98:
	movl	$.L.str.104, %eax
	retq
.LBB3_3:
	cmpl	$-1059, %edi                    # imm = 0xFBDD
	je	.LBB3_89
# %bb.4:
	cmpl	$-1058, %edi                    # imm = 0xFBDE
	jne	.LBB3_99
# %bb.5:
	movl	$.L.str.94, %eax
	retq
.LBB3_7:
	cmpl	$-1057, %edi                    # imm = 0xFBDF
	jne	.LBB3_99
# %bb.8:
	movl	$.L.str.93, %eax
	retq
.LBB3_93:
	movl	$.L.str.99, %eax
	retq
.LBB3_89:
	movl	$.L.str.95, %eax
	retq
.LBB3_87:
	movl	$.L.str.91, %eax
	retq
.LBB3_96:
	movl	$.L.str.102, %eax
	retq
.LBB3_95:
	movl	$.L.str.101, %eax
	retq
.LBB3_90:
	movl	$.L.str.96, %eax
	retq
.LBB3_97:
	movl	$.L.str.103, %eax
	retq
.LBB3_92:
	movl	$.L.str.98, %eax
	retq
.LBB3_91:
	movl	$.L.str.97, %eax
	retq
.LBB3_86:
	movl	$.L.str.90, %eax
	retq
.LBB3_20:
	movl	$.L.str.24, %eax
	retq
.LBB3_19:
	movl	$.L.str.23, %eax
	retq
.LBB3_39:
	movl	$.L.str.43, %eax
	retq
.LBB3_42:
	movl	$.L.str.46, %eax
	retq
.LBB3_36:
	movl	$.L.str.40, %eax
	retq
.LBB3_40:
	movl	$.L.str.44, %eax
	retq
.LBB3_22:
	movl	$.L.str.26, %eax
	retq
.LBB3_29:
	movl	$.L.str.33, %eax
	retq
.LBB3_25:
	movl	$.L.str.29, %eax
	retq
.LBB3_51:
	movl	$.L.str.55, %eax
	retq
.LBB3_44:
	movl	$.L.str.48, %eax
	retq
.LBB3_52:
	movl	$.L.str.56, %eax
	retq
.LBB3_31:
	movl	$.L.str.35, %eax
	retq
.LBB3_38:
	movl	$.L.str.42, %eax
	retq
.LBB3_60:
	movl	$.L.str.64, %eax
	retq
.LBB3_37:
	movl	$.L.str.41, %eax
	retq
.LBB3_13:
	movl	$.L.str.17, %eax
	retq
.LBB3_46:
	movl	$.L.str.50, %eax
	retq
.LBB3_53:
	movl	$.L.str.57, %eax
	retq
.LBB3_12:
	movl	$.L.str.16, %eax
	retq
.LBB3_54:
	movl	$.L.str.58, %eax
	retq
.LBB3_15:
	movl	$.L.str.19, %eax
	retq
.LBB3_69:
	movl	$.L.str.73, %eax
	retq
.LBB3_27:
	movl	$.L.str.31, %eax
	retq
.LBB3_67:
	movl	$.L.str.71, %eax
	retq
.LBB3_71:
	movl	$.L.str.75, %eax
	retq
.LBB3_62:
	movl	$.L.str.66, %eax
	retq
.LBB3_30:
	movl	$.L.str.34, %eax
	retq
.LBB3_33:
	movl	$.L.str.37, %eax
	retq
.LBB3_70:
	movl	$.L.str.74, %eax
	retq
.LBB3_21:
	movl	$.L.str.25, %eax
	retq
.LBB3_64:
	movl	$.L.str.68, %eax
	retq
.LBB3_72:
	movl	$.L.str.76, %eax
	retq
.LBB3_50:
	movl	$.L.str.54, %eax
	retq
.LBB3_23:
	movl	$.L.str.27, %eax
	retq
.LBB3_57:
	movl	$.L.str.61, %eax
	retq
.LBB3_32:
	movl	$.L.str.36, %eax
	retq
.LBB3_68:
	movl	$.L.str.72, %eax
	retq
.LBB3_63:
	movl	$.L.str.67, %eax
	retq
.LBB3_66:
	movl	$.L.str.70, %eax
	retq
.LBB3_56:
	movl	$.L.str.60, %eax
	retq
.LBB3_55:
	movl	$.L.str.59, %eax
	retq
.LBB3_65:
	movl	$.L.str.69, %eax
	retq
.LBB3_59:
	movl	$.L.str.63, %eax
	retq
.LBB3_58:
	movl	$.L.str.62, %eax
	retq
.LBB3_16:
	movl	$.L.str.20, %eax
	retq
.LBB3_61:
	movl	$.L.str.65, %eax
	retq
.LBB3_45:
	movl	$.L.str.49, %eax
	retq
.LBB3_49:
	movl	$.L.str.53, %eax
	retq
.LBB3_47:
	movl	$.L.str.51, %eax
	retq
.LBB3_43:
	movl	$.L.str.47, %eax
	retq
.LBB3_48:
	movl	$.L.str.52, %eax
	retq
.LBB3_14:
	movl	$.L.str.18, %eax
	retq
.LBB3_34:
	movl	$.L.str.38, %eax
	retq
.LBB3_41:
	movl	$.L.str.45, %eax
	retq
.LBB3_35:
	movl	$.L.str.39, %eax
	retq
.LBB3_26:
	movl	$.L.str.30, %eax
	retq
.LBB3_18:
	movl	$.L.str.22, %eax
	retq
.LBB3_28:
	movl	$.L.str.32, %eax
	retq
.LBB3_24:
	movl	$.L.str.28, %eax
	retq
.LBB3_17:
	movl	$.L.str.21, %eax
	retq
.LBB3_74:
	movl	$.L.str.78, %eax
	retq
.LBB3_81:
	movl	$.L.str.85, %eax
	retq
.LBB3_88:
	movl	$.L.str.92, %eax
	retq
.LBB3_76:
	movl	$.L.str.80, %eax
	retq
.LBB3_83:
	movl	$.L.str.87, %eax
	retq
.LBB3_82:
	movl	$.L.str.86, %eax
	retq
.LBB3_78:
	movl	$.L.str.82, %eax
	retq
.LBB3_84:
	movl	$.L.str.88, %eax
	retq
.LBB3_80:
	movl	$.L.str.84, %eax
	retq
.LBB3_94:
	movl	$.L.str.100, %eax
	retq
.LBB3_79:
	movl	$.L.str.83, %eax
	retq
.LBB3_75:
	movl	$.L.str.79, %eax
	retq
.LBB3_77:
	movl	$.L.str.81, %eax
	retq
.LBB3_99:
	movl	$.L.str.105, %eax
	retq
.Lfunc_end3:
	.size	_Z16clGetErrorStringi, .Lfunc_end3-_Z16clGetErrorStringi
	.cfi_endproc
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
.LJTI3_0:
	.quad	.LBB3_98
	.quad	.LBB3_97
	.quad	.LBB3_96
	.quad	.LBB3_95
	.quad	.LBB3_93
	.quad	.LBB3_92
	.quad	.LBB3_91
	.quad	.LBB3_90
	.quad	.LBB3_86
	.quad	.LBB3_87
.LJTI3_1:
	.quad	.LBB3_85
	.quad	.LBB3_84
	.quad	.LBB3_83
	.quad	.LBB3_82
	.quad	.LBB3_81
	.quad	.LBB3_80
	.quad	.LBB3_79
	.quad	.LBB3_78
	.quad	.LBB3_77
	.quad	.LBB3_76
	.quad	.LBB3_75
	.quad	.LBB3_74
	.quad	.LBB3_88
	.quad	.LBB3_94
.LJTI3_2:
	.quad	.LBB3_73
	.quad	.LBB3_72
	.quad	.LBB3_71
	.quad	.LBB3_70
	.quad	.LBB3_69
	.quad	.LBB3_68
	.quad	.LBB3_67
	.quad	.LBB3_66
	.quad	.LBB3_65
	.quad	.LBB3_64
	.quad	.LBB3_63
	.quad	.LBB3_62
	.quad	.LBB3_61
	.quad	.LBB3_60
	.quad	.LBB3_59
	.quad	.LBB3_58
	.quad	.LBB3_57
	.quad	.LBB3_56
	.quad	.LBB3_55
	.quad	.LBB3_54
	.quad	.LBB3_53
	.quad	.LBB3_52
	.quad	.LBB3_51
	.quad	.LBB3_50
	.quad	.LBB3_49
	.quad	.LBB3_48
	.quad	.LBB3_47
	.quad	.LBB3_46
	.quad	.LBB3_45
	.quad	.LBB3_44
	.quad	.LBB3_43
	.quad	.LBB3_42
	.quad	.LBB3_41
	.quad	.LBB3_40
	.quad	.LBB3_39
	.quad	.LBB3_38
	.quad	.LBB3_37
	.quad	.LBB3_36
	.quad	.LBB3_35
	.quad	.LBB3_34
	.quad	.LBB3_33
	.quad	.LBB3_32
	.quad	.LBB3_31
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_99
	.quad	.LBB3_30
	.quad	.LBB3_29
	.quad	.LBB3_28
	.quad	.LBB3_27
	.quad	.LBB3_26
	.quad	.LBB3_25
	.quad	.LBB3_24
	.quad	.LBB3_23
	.quad	.LBB3_99
	.quad	.LBB3_22
	.quad	.LBB3_21
	.quad	.LBB3_20
	.quad	.LBB3_19
	.quad	.LBB3_18
	.quad	.LBB3_17
	.quad	.LBB3_16
	.quad	.LBB3_15
	.quad	.LBB3_14
	.quad	.LBB3_13
	.quad	.LBB3_12
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"typedef struct{                                         \nfloat position[3];                                    \nfloat velocity[3];                                   \n}particle;                                                       \n__kernel                                                 \nvoid kernel_upload(                                       \n   const int iter,                                          \n   const int nump,                                        \n   __global particle* X)                                    \n{int index = get_global_id(0);                           \nif(index < nump) {                                                \n   X[index].velocity[0] = (3*index + iter) % nump;             \n   X[index].velocity[1] = (4*index + iter) % nump;            \n   X[index].velocity[2] = (5*index + iter) % nump;             \n   X[index].position[0] = X[index].position[0] + X[index].velocity[0];             \n   X[index].position[1] = X[index].position[1] + X[index].velocity[1];             \n   X[index].position[2] = X[index].position[2] + X[index].velocity[2];             \n}                                                          \n}                                                        \n"
	.size	.L.str, 1217

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

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

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Build error: %s\n"
	.size	.L.str.5, 17

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"kernel_upload"
	.size	.L.str.6, 14

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Error en %d i en dim %d, resultado %f i %f\n"
	.size	.L.str.8, 44

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"NUM_PARTICLES: %d\n"
	.size	.L.str.11, 19

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"BLOCK_SIZE: %d\n"
	.size	.L.str.12, 16

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

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"Execution time on CPU: %f\n"
	.size	.L.str.14, 27

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"Execution time on GPU: %f\n"
	.size	.L.str.15, 27

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"CL_SUCCESS"
	.size	.L.str.16, 11

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"CL_DEVICE_NOT_FOUND"
	.size	.L.str.17, 20

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"CL_DEVICE_NOT_AVAILABLE"
	.size	.L.str.18, 24

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"CL_COMPILER_NOT_AVAILABLE"
	.size	.L.str.19, 26

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"CL_MEM_OBJECT_ALLOCATION_FAILURE"
	.size	.L.str.20, 33

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"CL_OUT_OF_RESOURCES"
	.size	.L.str.21, 20

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"CL_OUT_OF_HOST_MEMORY"
	.size	.L.str.22, 22

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"CL_PROFILING_INFO_NOT_AVAILABLE"
	.size	.L.str.23, 32

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	"CL_MEM_COPY_OVERLAP"
	.size	.L.str.24, 20

	.type	.L.str.25,@object               # @.str.25
.L.str.25:
	.asciz	"CL_IMAGE_FORMAT_MISMATCH"
	.size	.L.str.25, 25

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"CL_IMAGE_FORMAT_NOT_SUPPORTED"
	.size	.L.str.26, 30

	.type	.L.str.27,@object               # @.str.27
.L.str.27:
	.asciz	"CL_MAP_FAILURE"
	.size	.L.str.27, 15

	.type	.L.str.28,@object               # @.str.28
.L.str.28:
	.asciz	"CL_MISALIGNED_SUB_BUFFER_OFFSET"
	.size	.L.str.28, 32

	.type	.L.str.29,@object               # @.str.29
.L.str.29:
	.asciz	"CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"
	.size	.L.str.29, 45

	.type	.L.str.30,@object               # @.str.30
.L.str.30:
	.asciz	"CL_COMPILE_PROGRAM_FAILURE"
	.size	.L.str.30, 27

	.type	.L.str.31,@object               # @.str.31
.L.str.31:
	.asciz	"CL_LINKER_NOT_AVAILABLE"
	.size	.L.str.31, 24

	.type	.L.str.32,@object               # @.str.32
.L.str.32:
	.asciz	"CL_LINK_PROGRAM_FAILURE"
	.size	.L.str.32, 24

	.type	.L.str.33,@object               # @.str.33
.L.str.33:
	.asciz	"CL_DEVICE_PARTITION_FAILED"
	.size	.L.str.33, 27

	.type	.L.str.34,@object               # @.str.34
.L.str.34:
	.asciz	"CL_KERNEL_ARG_INFO_NOT_AVAILABLE"
	.size	.L.str.34, 33

	.type	.L.str.35,@object               # @.str.35
.L.str.35:
	.asciz	"CL_INVALID_VALUE"
	.size	.L.str.35, 17

	.type	.L.str.36,@object               # @.str.36
.L.str.36:
	.asciz	"CL_INVALID_DEVICE_TYPE"
	.size	.L.str.36, 23

	.type	.L.str.37,@object               # @.str.37
.L.str.37:
	.asciz	"CL_INVALID_PLATFORM"
	.size	.L.str.37, 20

	.type	.L.str.38,@object               # @.str.38
.L.str.38:
	.asciz	"CL_INVALID_DEVICE"
	.size	.L.str.38, 18

	.type	.L.str.39,@object               # @.str.39
.L.str.39:
	.asciz	"CL_INVALID_CONTEXT"
	.size	.L.str.39, 19

	.type	.L.str.40,@object               # @.str.40
.L.str.40:
	.asciz	"CL_INVALID_QUEUE_PROPERTIES"
	.size	.L.str.40, 28

	.type	.L.str.41,@object               # @.str.41
.L.str.41:
	.asciz	"CL_INVALID_COMMAND_QUEUE"
	.size	.L.str.41, 25

	.type	.L.str.42,@object               # @.str.42
.L.str.42:
	.asciz	"CL_INVALID_HOST_PTR"
	.size	.L.str.42, 20

	.type	.L.str.43,@object               # @.str.43
.L.str.43:
	.asciz	"CL_INVALID_MEM_OBJECT"
	.size	.L.str.43, 22

	.type	.L.str.44,@object               # @.str.44
.L.str.44:
	.asciz	"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"
	.size	.L.str.44, 35

	.type	.L.str.45,@object               # @.str.45
.L.str.45:
	.asciz	"CL_INVALID_IMAGE_SIZE"
	.size	.L.str.45, 22

	.type	.L.str.46,@object               # @.str.46
.L.str.46:
	.asciz	"CL_INVALID_SAMPLER"
	.size	.L.str.46, 19

	.type	.L.str.47,@object               # @.str.47
.L.str.47:
	.asciz	"CL_INVALID_BINARY"
	.size	.L.str.47, 18

	.type	.L.str.48,@object               # @.str.48
.L.str.48:
	.asciz	"CL_INVALID_BUILD_OPTIONS"
	.size	.L.str.48, 25

	.type	.L.str.49,@object               # @.str.49
.L.str.49:
	.asciz	"CL_INVALID_PROGRAM"
	.size	.L.str.49, 19

	.type	.L.str.50,@object               # @.str.50
.L.str.50:
	.asciz	"CL_INVALID_PROGRAM_EXECUTABLE"
	.size	.L.str.50, 30

	.type	.L.str.51,@object               # @.str.51
.L.str.51:
	.asciz	"CL_INVALID_KERNEL_NAME"
	.size	.L.str.51, 23

	.type	.L.str.52,@object               # @.str.52
.L.str.52:
	.asciz	"CL_INVALID_KERNEL_DEFINITION"
	.size	.L.str.52, 29

	.type	.L.str.53,@object               # @.str.53
.L.str.53:
	.asciz	"CL_INVALID_KERNEL"
	.size	.L.str.53, 18

	.type	.L.str.54,@object               # @.str.54
.L.str.54:
	.asciz	"CL_INVALID_ARG_INDEX"
	.size	.L.str.54, 21

	.type	.L.str.55,@object               # @.str.55
.L.str.55:
	.asciz	"CL_INVALID_ARG_VALUE"
	.size	.L.str.55, 21

	.type	.L.str.56,@object               # @.str.56
.L.str.56:
	.asciz	"CL_INVALID_ARG_SIZE"
	.size	.L.str.56, 20

	.type	.L.str.57,@object               # @.str.57
.L.str.57:
	.asciz	"CL_INVALID_KERNEL_ARGS"
	.size	.L.str.57, 23

	.type	.L.str.58,@object               # @.str.58
.L.str.58:
	.asciz	"CL_INVALID_WORK_DIMENSION"
	.size	.L.str.58, 26

	.type	.L.str.59,@object               # @.str.59
.L.str.59:
	.asciz	"CL_INVALID_WORK_GROUP_SIZE"
	.size	.L.str.59, 27

	.type	.L.str.60,@object               # @.str.60
.L.str.60:
	.asciz	"CL_INVALID_WORK_ITEM_SIZE"
	.size	.L.str.60, 26

	.type	.L.str.61,@object               # @.str.61
.L.str.61:
	.asciz	"CL_INVALID_GLOBAL_OFFSET"
	.size	.L.str.61, 25

	.type	.L.str.62,@object               # @.str.62
.L.str.62:
	.asciz	"CL_INVALID_EVENT_WAIT_LIST"
	.size	.L.str.62, 27

	.type	.L.str.63,@object               # @.str.63
.L.str.63:
	.asciz	"CL_INVALID_EVENT"
	.size	.L.str.63, 17

	.type	.L.str.64,@object               # @.str.64
.L.str.64:
	.asciz	"CL_INVALID_OPERATION"
	.size	.L.str.64, 21

	.type	.L.str.65,@object               # @.str.65
.L.str.65:
	.asciz	"CL_INVALID_GL_OBJECT"
	.size	.L.str.65, 21

	.type	.L.str.66,@object               # @.str.66
.L.str.66:
	.asciz	"CL_INVALID_BUFFER_SIZE"
	.size	.L.str.66, 23

	.type	.L.str.67,@object               # @.str.67
.L.str.67:
	.asciz	"CL_INVALID_MIP_LEVEL"
	.size	.L.str.67, 21

	.type	.L.str.68,@object               # @.str.68
.L.str.68:
	.asciz	"CL_INVALID_GLOBAL_WORK_SIZE"
	.size	.L.str.68, 28

	.type	.L.str.69,@object               # @.str.69
.L.str.69:
	.asciz	"CL_INVALID_PROPERTY"
	.size	.L.str.69, 20

	.type	.L.str.70,@object               # @.str.70
.L.str.70:
	.asciz	"CL_INVALID_IMAGE_DESCRIPTOR"
	.size	.L.str.70, 28

	.type	.L.str.71,@object               # @.str.71
.L.str.71:
	.asciz	"CL_INVALID_COMPILER_OPTIONS"
	.size	.L.str.71, 28

	.type	.L.str.72,@object               # @.str.72
.L.str.72:
	.asciz	"CL_INVALID_LINKER_OPTIONS"
	.size	.L.str.72, 26

	.type	.L.str.73,@object               # @.str.73
.L.str.73:
	.asciz	"CL_INVALID_DEVICE_PARTITION_COUNT"
	.size	.L.str.73, 34

	.type	.L.str.74,@object               # @.str.74
.L.str.74:
	.asciz	"CL_INVALID_PIPE_SIZE"
	.size	.L.str.74, 21

	.type	.L.str.75,@object               # @.str.75
.L.str.75:
	.asciz	"CL_INVALID_DEVICE_QUEUE"
	.size	.L.str.75, 24

	.type	.L.str.76,@object               # @.str.76
.L.str.76:
	.asciz	"CL_INVALID_SPEC_ID"
	.size	.L.str.76, 19

	.type	.L.str.77,@object               # @.str.77
.L.str.77:
	.asciz	"CL_MAX_SIZE_RESTRICTION_EXCEEDED"
	.size	.L.str.77, 33

	.type	.L.str.78,@object               # @.str.78
.L.str.78:
	.asciz	"CL_INVALID_D3D10_DEVICE_KHR"
	.size	.L.str.78, 28

	.type	.L.str.79,@object               # @.str.79
.L.str.79:
	.asciz	"CL_INVALID_D3D10_RESOURCE_KHR"
	.size	.L.str.79, 30

	.type	.L.str.80,@object               # @.str.80
.L.str.80:
	.asciz	"CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR"
	.size	.L.str.80, 39

	.type	.L.str.81,@object               # @.str.81
.L.str.81:
	.asciz	"CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR"
	.size	.L.str.81, 35

	.type	.L.str.82,@object               # @.str.82
.L.str.82:
	.asciz	"CL_INVALID_D3D11_DEVICE_KHR"
	.size	.L.str.82, 28

	.type	.L.str.83,@object               # @.str.83
.L.str.83:
	.asciz	"CL_INVALID_D3D11_RESOURCE_KHR"
	.size	.L.str.83, 30

	.type	.L.str.84,@object               # @.str.84
.L.str.84:
	.asciz	"CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR"
	.size	.L.str.84, 39

	.type	.L.str.85,@object               # @.str.85
.L.str.85:
	.asciz	"CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR"
	.size	.L.str.85, 35

	.type	.L.str.86,@object               # @.str.86
.L.str.86:
	.asciz	"CL_INVALID_DX9_MEDIA_ADAPTER_KHR"
	.size	.L.str.86, 33

	.type	.L.str.87,@object               # @.str.87
.L.str.87:
	.asciz	"CL_INVALID_DX9_MEDIA_SURFACE_KHR"
	.size	.L.str.87, 33

	.type	.L.str.88,@object               # @.str.88
.L.str.88:
	.asciz	"CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR"
	.size	.L.str.88, 42

	.type	.L.str.89,@object               # @.str.89
.L.str.89:
	.asciz	"CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR"
	.size	.L.str.89, 38

	.type	.L.str.90,@object               # @.str.90
.L.str.90:
	.asciz	"CL_INVALID_EGL_OBJECT_KHR"
	.size	.L.str.90, 26

	.type	.L.str.91,@object               # @.str.91
.L.str.91:
	.asciz	"CL_EGL_RESOURCE_NOT_ACQUIRED_KHR"
	.size	.L.str.91, 33

	.type	.L.str.92,@object               # @.str.92
.L.str.92:
	.asciz	"CL_PLATFORM_NOT_FOUND_KHR"
	.size	.L.str.92, 26

	.type	.L.str.93,@object               # @.str.93
.L.str.93:
	.asciz	"CL_DEVICE_PARTITION_FAILED_EXT"
	.size	.L.str.93, 31

	.type	.L.str.94,@object               # @.str.94
.L.str.94:
	.asciz	"CL_INVALID_PARTITION_COUNT_EXT"
	.size	.L.str.94, 31

	.type	.L.str.95,@object               # @.str.95
.L.str.95:
	.asciz	"CL_INVALID_PARTITION_NAME_EXT"
	.size	.L.str.95, 30

	.type	.L.str.96,@object               # @.str.96
.L.str.96:
	.asciz	"CL_INVALID_ACCELERATOR_INTEL"
	.size	.L.str.96, 29

	.type	.L.str.97,@object               # @.str.97
.L.str.97:
	.asciz	"CL_INVALID_ACCELERATOR_TYPE_INTEL"
	.size	.L.str.97, 34

	.type	.L.str.98,@object               # @.str.98
.L.str.98:
	.asciz	"CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL"
	.size	.L.str.98, 40

	.type	.L.str.99,@object               # @.str.99
.L.str.99:
	.asciz	"CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL"
	.size	.L.str.99, 40

	.type	.L.str.100,@object              # @.str.100
.L.str.100:
	.asciz	"CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR"
	.size	.L.str.100, 39

	.type	.L.str.101,@object              # @.str.101
.L.str.101:
	.asciz	"CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL"
	.size	.L.str.101, 38

	.type	.L.str.102,@object              # @.str.102
.L.str.102:
	.asciz	"CL_INVALID_VA_API_MEDIA_SURFACE_INTEL"
	.size	.L.str.102, 38

	.type	.L.str.103,@object              # @.str.103
.L.str.103:
	.asciz	"CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL"
	.size	.L.str.103, 47

	.type	.L.str.104,@object              # @.str.104
.L.str.104:
	.asciz	"CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL"
	.size	.L.str.104, 43

	.type	.L.str.105,@object              # @.str.105
.L.str.105:
	.asciz	"CL_UNKNOWN_ERROR"
	.size	.L.str.105, 17

	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Error: Failed to create a compute context!"
	.size	.Lstr, 43

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Error: Failed to create a command queue!"
	.size	.Lstr.1, 41

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"Computing SAXPY on the CPU... Done"
	.size	.Lstr.2, 35

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"Computing SACY on the GPU... Done"
	.size	.Lstr.3, 34

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"Comparing the output for each implementation... Incorrect!"
	.size	.Lstr.4, 59

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"Comparing the output for each implementation... Correct!"
	.size	.Lstr.5, 57

	.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 mykernel
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
