1*8468e674SAndrey Tretyakov; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV 2*8468e674SAndrey Tretyakov 3*8468e674SAndrey Tretyakov; CHECK-SPIRV: OpFNegate 4*8468e674SAndrey Tretyakov; CHECK-SPIRV: OpFNegate 5*8468e674SAndrey Tretyakov; CHECK-SPIRV: OpFNegate 6*8468e674SAndrey Tretyakov; CHECK-SPIRV: OpFNegate 7*8468e674SAndrey Tretyakov 8*8468e674SAndrey Tretyakov;; #pragma OPENCL EXTENSION cl_khr_fp64 : enable 9*8468e674SAndrey Tretyakov;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable 10*8468e674SAndrey Tretyakov;; 11*8468e674SAndrey Tretyakov;; __kernel void foo(double a1, __global half *h, __global float *b0, __global double *b1, __global double8 *d) { 12*8468e674SAndrey Tretyakov;; *h = -*h; 13*8468e674SAndrey Tretyakov;; *b0 = -*b0; 14*8468e674SAndrey Tretyakov;; *b1 = -a1; 15*8468e674SAndrey Tretyakov;; *d = -*d; 16*8468e674SAndrey Tretyakov;; } 17*8468e674SAndrey Tretyakov 18*8468e674SAndrey Tretyakovdefine dso_local spir_kernel void @foo(double noundef %a1, half addrspace(1)* noundef %h, float addrspace(1)* noundef %b0, double addrspace(1)* noundef %b1, <8 x double> addrspace(1)* noundef %d) { 19*8468e674SAndrey Tretyakoventry: 20*8468e674SAndrey Tretyakov %a1.addr = alloca double, align 8 21*8468e674SAndrey Tretyakov %h.addr = alloca half addrspace(1)*, align 4 22*8468e674SAndrey Tretyakov %b0.addr = alloca float addrspace(1)*, align 4 23*8468e674SAndrey Tretyakov %b1.addr = alloca double addrspace(1)*, align 4 24*8468e674SAndrey Tretyakov %d.addr = alloca <8 x double> addrspace(1)*, align 4 25*8468e674SAndrey Tretyakov store double %a1, double* %a1.addr, align 8 26*8468e674SAndrey Tretyakov store half addrspace(1)* %h, half addrspace(1)** %h.addr, align 4 27*8468e674SAndrey Tretyakov store float addrspace(1)* %b0, float addrspace(1)** %b0.addr, align 4 28*8468e674SAndrey Tretyakov store double addrspace(1)* %b1, double addrspace(1)** %b1.addr, align 4 29*8468e674SAndrey Tretyakov store <8 x double> addrspace(1)* %d, <8 x double> addrspace(1)** %d.addr, align 4 30*8468e674SAndrey Tretyakov %0 = load half addrspace(1)*, half addrspace(1)** %h.addr, align 4 31*8468e674SAndrey Tretyakov %1 = load half, half addrspace(1)* %0, align 2 32*8468e674SAndrey Tretyakov %fneg = fneg half %1 33*8468e674SAndrey Tretyakov %2 = load half addrspace(1)*, half addrspace(1)** %h.addr, align 4 34*8468e674SAndrey Tretyakov store half %fneg, half addrspace(1)* %2, align 2 35*8468e674SAndrey Tretyakov %3 = load float addrspace(1)*, float addrspace(1)** %b0.addr, align 4 36*8468e674SAndrey Tretyakov %4 = load float, float addrspace(1)* %3, align 4 37*8468e674SAndrey Tretyakov %fneg1 = fneg float %4 38*8468e674SAndrey Tretyakov %5 = load float addrspace(1)*, float addrspace(1)** %b0.addr, align 4 39*8468e674SAndrey Tretyakov store float %fneg1, float addrspace(1)* %5, align 4 40*8468e674SAndrey Tretyakov %6 = load double, double* %a1.addr, align 8 41*8468e674SAndrey Tretyakov %fneg2 = fneg double %6 42*8468e674SAndrey Tretyakov %7 = load double addrspace(1)*, double addrspace(1)** %b1.addr, align 4 43*8468e674SAndrey Tretyakov store double %fneg2, double addrspace(1)* %7, align 8 44*8468e674SAndrey Tretyakov %8 = load <8 x double> addrspace(1)*, <8 x double> addrspace(1)** %d.addr, align 4 45*8468e674SAndrey Tretyakov %9 = load <8 x double>, <8 x double> addrspace(1)* %8, align 64 46*8468e674SAndrey Tretyakov %fneg3 = fneg <8 x double> %9 47*8468e674SAndrey Tretyakov %10 = load <8 x double> addrspace(1)*, <8 x double> addrspace(1)** %d.addr, align 4 48*8468e674SAndrey Tretyakov store <8 x double> %fneg3, <8 x double> addrspace(1)* %10, align 64 49*8468e674SAndrey Tretyakov ret void 50*8468e674SAndrey Tretyakov} 51