1;; __kernel void test( __global float4 *p, __global half *f ) 2;; { 3;; __private float4 data; 4;; data = p[0]; 5;; vstorea_half4_rtp( data, 0, f ); 6;; } 7 8; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s 9; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} 10 11; CHECK-DAG: OpCapability Float16Buffer 12; CHECK-DAG: OpCapability Float16 13 14define spir_kernel void @test(<4 x float> addrspace(1)* %p, half addrspace(1)* %f) { 15entry: 16 %p.addr = alloca <4 x float> addrspace(1)*, align 8 17 %f.addr = alloca half addrspace(1)*, align 8 18 %data = alloca <4 x float>, align 16 19 store <4 x float> addrspace(1)* %p, <4 x float> addrspace(1)** %p.addr, align 8 20 store half addrspace(1)* %f, half addrspace(1)** %f.addr, align 8 21 %0 = load <4 x float> addrspace(1)*, <4 x float> addrspace(1)** %p.addr, align 8 22 %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %0, i64 0 23 %1 = load <4 x float>, <4 x float> addrspace(1)* %arrayidx, align 16 24 store <4 x float> %1, <4 x float>* %data, align 16 25 %2 = load <4 x float>, <4 x float>* %data, align 16 26 %3 = load half addrspace(1)*, half addrspace(1)** %f.addr, align 8 27 call spir_func void @_Z17vstorea_half4_rtpDv4_fmPU3AS1Dh(<4 x float> %2, i64 0, half addrspace(1)* %3) 28 ret void 29} 30 31declare spir_func void @_Z17vstorea_half4_rtpDv4_fmPU3AS1Dh(<4 x float>, i64, half addrspace(1)*) 32