This is an archive of the discontinued LLVM Phabricator instance.

Differential D142061

[SPIR-V] Add -opaque-pointers=0 to some LIT tests
ClosedPublic

Authored by mpaszkowski on Jan 18 2023, 3:15 PM.

Details

Reviewers
nikic
zuban32
konrad.trifunovic
iliya-diyachkov
Commits
rG786cb151d938: [SPIR-V] Add -opaque-pointers=0 to some LIT tests
Summary

@nikic in D141912 removed auto-detection of typed pointers and added -opaque-pointers=0 to LITs which still require changes in other backends. Unfortunately, we still have such tests in the SPIR-V backend. This patch adds -opaque-pointers=0 to the 50 LIT tests which still require changes. I will create a tracker issue on GitHub for those tests and convert as many as possible to opaque pointers starting next week. However, some will need to wait for Joshua's changes introducing target extension types for OpenCL.

@nikic Please accept the patch and keep us in the loop. I will include you in the upcoming patches.

Diff Detail

Event Timeline

mpaszkowski created this revision.Jan 18 2023, 3:15 PM
Herald added a project: Restricted Project. · View Herald TranscriptJan 18 2023, 3:15 PM
Herald added subscribers: StephenFan, wenlei, ThomasRaoux, Anastasia. · View Herald Transcript
mpaszkowski requested review of this revision.Jan 18 2023, 3:15 PM
Herald added a project: Restricted Project. · View Herald TranscriptJan 18 2023, 3:15 PM
Herald added a subscriber: llvm-commits. · View Herald Transcript
mpaszkowski added a comment.Jan 18 2023, 3:20 PM

@nikic I have created the issue for tracking the SPIR-V backend LIT tests with -opaque-pointers=0 here: https://github.com/llvm/llvm-project/issues/60133

Harbormaster completed remote builds in B208612: Diff 490312.Jan 18 2023, 4:24 PM
nikic accepted this revision.Jan 19 2023, 2:08 AM

LGTM

This revision is now accepted and ready to land.Jan 19 2023, 2:08 AM
This revision was landed with ongoing or failed builds.Jan 19 2023, 5:02 AM
Closed by commit rG786cb151d938: [SPIR-V] Add -opaque-pointers=0 to some LIT tests (authored by mpaszkowski). · Explain Why
This revision was automatically updated to reflect the committed changes.
mpaszkowski added a commit: rG786cb151d938: [SPIR-V] Add -opaque-pointers=0 to some LIT tests.

Revision Contents

Diff 490462

llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
;; This test checks that Invoke parameter of OpEnueueKernel instruction meet the ;; This test checks that Invoke parameter of OpEnueueKernel instruction meet the
;; following specification requirements in case of enqueueing empty block: ;; following specification requirements in case of enqueueing empty block:
;; "Invoke must be an OpFunction whose OpTypeFunction operand has: ;; "Invoke must be an OpFunction whose OpTypeFunction operand has:
;; - Result Type must be OpTypeVoid. ;; - Result Type must be OpTypeVoid.
;; - The first parameter must have a type of OpTypePointer to an 8-bit OpTypeInt. ;; - The first parameter must have a type of OpTypePointer to an 8-bit OpTypeInt.
;; - An optional list of parameters, each of which must have a type of OpTypePointer to the Workgroup Storage Class. ;; - An optional list of parameters, each of which must have a type of OpTypePointer to the Workgroup Storage Class.
;; ... " ;; ... "
▲ Show 20 Lines • Show All 55 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/SampledImageRetType.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
%opencl.image1d_ro_t = type opaque %opencl.image1d_ro_t = type opaque
; CHECK: %[[#image1d_t:]] = OpTypeImage ; CHECK: %[[#image1d_t:]] = OpTypeImage
%opencl.sampler_t = type opaque %opencl.sampler_t = type opaque
; CHECK: %[[#sampler_t:]] = OpTypeSampler ; CHECK: %[[#sampler_t:]] = OpTypeSampler
; CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage ; CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage
declare dso_local spir_func i8 addrspace(4)* @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) local_unnamed_addr declare dso_local spir_func i8 addrspace(4)* @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) local_unnamed_addr
Show All 20 Lines

llvm/test/CodeGen/SPIRV/atomicrmw.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
; CHECK: %[[#Int:]] = OpTypeInt 32 0 ; CHECK: %[[#Int:]] = OpTypeInt 32 0
; CHECK-DAG: %[[#Scope_Device:]] = OpConstant %[[#Int]] 1 {{$}} ; CHECK-DAG: %[[#Scope_Device:]] = OpConstant %[[#Int]] 1 {{$}}
; CHECK-DAG: %[[#MemSem_Relaxed:]] = OpConstant %[[#Int]] 0 ; CHECK-DAG: %[[#MemSem_Relaxed:]] = OpConstant %[[#Int]] 0
; CHECK-DAG: %[[#MemSem_Acquire:]] = OpConstant %[[#Int]] 2 ; CHECK-DAG: %[[#MemSem_Acquire:]] = OpConstant %[[#Int]] 2
; CHECK-DAG: %[[#MemSem_Release:]] = OpConstant %[[#Int]] 4 {{$}} ; CHECK-DAG: %[[#MemSem_Release:]] = OpConstant %[[#Int]] 4 {{$}}
; CHECK-DAG: %[[#MemSem_AcquireRelease:]] = OpConstant %[[#Int]] 8 ; CHECK-DAG: %[[#MemSem_AcquireRelease:]] = OpConstant %[[#Int]] 8
; CHECK-DAG: %[[#MemSem_SequentiallyConsistent:]] = OpConstant %[[#Int]] 16 ; CHECK-DAG: %[[#MemSem_SequentiallyConsistent:]] = OpConstant %[[#Int]] 16
▲ Show 20 Lines • Show All 48 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/constant/global-constants.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
@global = addrspace(1) constant i32 1 ; OpenCL global memory @global = addrspace(1) constant i32 1 ; OpenCL global memory
@constant = addrspace(2) constant i32 2 ; OpenCL constant memory @constant = addrspace(2) constant i32 2 ; OpenCL constant memory
@local = addrspace(3) constant i32 3 ; OpenCL local memory @local = addrspace(3) constant i32 3 ; OpenCL local memory
define i32 @getGlobal1() { define i32 @getGlobal1() {
%g = load i32, i32 addrspace(1)* @global %g = load i32, i32 addrspace(1)* @global
ret i32 %g ret i32 %g
Show All 29 Lines

llvm/test/CodeGen/SPIRV/constant/local-null-constants.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
;; OpenCL global memory ;; OpenCL global memory
define i32 addrspace(1)* @getConstant1() { define i32 addrspace(1)* @getConstant1() {
ret i32 addrspace(1)* null ret i32 addrspace(1)* null
} }
;; OpenCL constant memory ;; OpenCL constant memory
define i32 addrspace(2)* @getConstant2() { define i32 addrspace(2)* @getConstant2() {
Show All 18 Lines

llvm/test/CodeGen/SPIRV/function/alloca-load-store.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; CHECK-DAG: OpName %[[#BAR:]] "bar" ; CHECK-DAG: OpName %[[#BAR:]] "bar"
; CHECK-DAG: OpName %[[#FOO:]] "foo" ; CHECK-DAG: OpName %[[#FOO:]] "foo"
; CHECK-DAG: OpName %[[#GOO:]] "goo" ; CHECK-DAG: OpName %[[#GOO:]] "goo"
; CHECK: %[[#INT:]] = OpTypeInt 32 ; CHECK: %[[#INT:]] = OpTypeInt 32
; CHECK-DAG: %[[#STACK_PTR:]] = OpTypePointer Function %[[#INT]] ; CHECK-DAG: %[[#STACK_PTR:]] = OpTypePointer Function %[[#INT]]
; CHECK-DAG: %[[#GLOBAL_PTR:]] = OpTypePointer CrossWorkgroup %[[#INT]] ; CHECK-DAG: %[[#GLOBAL_PTR:]] = OpTypePointer CrossWorkgroup %[[#INT]]
▲ Show 20 Lines • Show All 52 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/half_extension.ll

;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable ;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable
;; half test() ;; half test()
;; { ;; {
;; half x = 0.1f; ;; half x = 0.1f;
;; x += 2.0f; ;; x += 2.0f;
;; half y = x + x; ;; half y = x + x;
;; return y; ;; return y;
;; } ;; }
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: OpCapability Float16Buffer ; CHECK-SPIRV-DAG: OpCapability Float16Buffer
; CHECK-SPIRV-DAG: OpCapability Float16 ; CHECK-SPIRV-DAG: OpCapability Float16
define spir_func half @test() { define spir_func half @test() {
entry: entry:
%x = alloca half, align 2 %x = alloca half, align 2
%y = alloca half, align 2 %y = alloca half, align 2
Show All 13 Lines

llvm/test/CodeGen/SPIRV/half_no_extension.ll

;; __kernel void test( __global float4 *p, __global half *f ) ;; __kernel void test( __global float4 *p, __global half *f )
;; { ;; {
;; __private float4 data; ;; __private float4 data;
;; data = p[0]; ;; data = p[0];
;; vstorea_half4_rtp( data, 0, f ); ;; vstorea_half4_rtp( data, 0, f );
;; } ;; }
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: OpCapability Float16Buffer ; CHECK-SPIRV: OpCapability Float16Buffer
; CHECK-SPIRV-NOT: OpCapability Float16 ; CHECK-SPIRV-NOT: OpCapability Float16
define spir_kernel void @test(<4 x float> addrspace(1)* %p, half addrspace(1)* %f) { define spir_kernel void @test(<4 x float> addrspace(1)* %p, half addrspace(1)* %f) {
entry: entry:
%p.addr = alloca <4 x float> addrspace(1)*, align 8 %p.addr = alloca <4 x float> addrspace(1)*, align 8
%f.addr = alloca half addrspace(1)*, align 8 %f.addr = alloca half addrspace(1)*, align 8
Show All 14 Lines

llvm/test/CodeGen/SPIRV/image-unoptimized.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; CHECK: %[[#TypeImage:]] = OpTypeImage ; CHECK: %[[#TypeImage:]] = OpTypeImage
; CHECK: %[[#TypeSampler:]] = OpTypeSampler ; CHECK: %[[#TypeSampler:]] = OpTypeSampler
; CHECK-DAG: %[[#TypeImagePtr:]] = OpTypePointer {{.*}} %[[#TypeImage]] ; CHECK-DAG: %[[#TypeImagePtr:]] = OpTypePointer {{.*}} %[[#TypeImage]]
; CHECK-DAG: %[[#TypeSamplerPtr:]] = OpTypePointer {{.*}} %[[#TypeSampler]] ; CHECK-DAG: %[[#TypeSamplerPtr:]] = OpTypePointer {{.*}} %[[#TypeSampler]]
; CHECK: %[[#srcimg:]] = OpFunctionParameter %[[#TypeImage]] ; CHECK: %[[#srcimg:]] = OpFunctionParameter %[[#TypeImage]]
; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#TypeSampler]] ; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#TypeSampler]]
▲ Show 20 Lines • Show All 80 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/image.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#VOID_TY:]] = OpTypeVoid ; CHECK-SPIRV: %[[#VOID_TY:]] = OpTypeVoid
; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly ; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly
; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly ; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly
; CHECK-SPIRV-NOT: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly ; CHECK-SPIRV-NOT: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly
; CHECK-SPIRV: OpImageSampleExplicitLod ; CHECK-SPIRV: OpImageSampleExplicitLod
; CHECK-SPIRV: OpImageWrite ; CHECK-SPIRV: OpImageWrite
Show All 22 Lines

llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#TypeImage:]] = OpTypeImage ; CHECK-SPIRV: %[[#TypeImage:]] = OpTypeImage
; CHECK-SPIRV-NOT: OpTypeImage ; CHECK-SPIRV-NOT: OpTypeImage
; CHECK-SPIRV: %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]] ; CHECK-SPIRV: %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]]
; CHECK-SPIRV: %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]] ; CHECK-SPIRV: %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]]
; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#TypeImage]] ; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#TypeImage]]
; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#TypeImage]] ; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#TypeImage]]
; CHECK-SPIRV: %[[#ParamID:]] = OpFunctionParameter %[[#TypeImage]] ; CHECK-SPIRV: %[[#ParamID:]] = OpFunctionParameter %[[#TypeImage]]
Show All 19 Lines

llvm/test/CodeGen/SPIRV/image_dim.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: OpCapability Sampled1D ; CHECK-SPIRV-DAG: OpCapability Sampled1D
; CHECK-SPIRV-DAG: OpCapability SampledBuffer ; CHECK-SPIRV-DAG: OpCapability SampledBuffer
%opencl.image1d_t = type opaque %opencl.image1d_t = type opaque
%opencl.image1d_buffer_t = type opaque %opencl.image1d_buffer_t = type opaque
define spir_kernel void @image_d(%opencl.image1d_t addrspace(1)* %image1d_td6, %opencl.image1d_buffer_t addrspace(1)* %image1d_buffer_td8) { define spir_kernel void @image_d(%opencl.image1d_t addrspace(1)* %image1d_td6, %opencl.image1d_buffer_t addrspace(1)* %image1d_buffer_td8) {
entry: entry:
ret void ret void
} }

llvm/test/CodeGen/SPIRV/image_store.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
;; Image types may be represented in two ways while translating to SPIR-V: ;; Image types may be represented in two ways while translating to SPIR-V:
;; - OpenCL form, for example, '%opencl.image2d_ro_t', ;; - OpenCL form, for example, '%opencl.image2d_ro_t',
;; - SPIR-V form, for example, '%spirv.Image._void_1_0_0_0_0_0_0', ;; - SPIR-V form, for example, '%spirv.Image._void_1_0_0_0_0_0_0',
;; but it is still one type which should be translated to one SPIR-V type. ;; but it is still one type which should be translated to one SPIR-V type.
;; ;;
;; The test checks that the code below is successfully translated and only one ;; The test checks that the code below is successfully translated and only one
;; SPIR-V type for images is generated. ;; SPIR-V type for images is generated.
Show All 14 Lines

llvm/test/CodeGen/SPIRV/link-attribute.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
%opencl.image2d_t = type opaque %opencl.image2d_t = type opaque
; CHECK: OpDecorate %[[#ID:]] LinkageAttributes "imageSampler" Export ; CHECK: OpDecorate %[[#ID:]] LinkageAttributes "imageSampler" Export
; CHECK: %[[#ID]] = OpVariable %[[#]] UniformConstant %[[#]] ; CHECK: %[[#ID]] = OpVariable %[[#]] UniformConstant %[[#]]
@imageSampler = addrspace(2) constant i32 36, align 4 @imageSampler = addrspace(2) constant i32 36, align 4
Show All 24 Lines

llvm/test/CodeGen/SPIRV/linked-list.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
%struct.Node = type { %struct.Node.0 addrspace(1)* } %struct.Node = type { %struct.Node.0 addrspace(1)* }
; CHECK: %[[#]] = OpTypeOpaque "struct.Node.0" ; CHECK: %[[#]] = OpTypeOpaque "struct.Node.0"
%struct.Node.0 = type opaque %struct.Node.0 = type opaque
define spir_kernel void @create_linked_lists(%struct.Node addrspace(1)* nocapture %pNodes, i32 addrspace(1)* nocapture %allocation_index, i32 %list_length) { define spir_kernel void @create_linked_lists(%struct.Node addrspace(1)* nocapture %pNodes, i32 addrspace(1)* nocapture %allocation_index, i32 %list_length) {
entry: entry:
ret void ret void
} }

llvm/test/CodeGen/SPIRV/literal-struct.ll

;; This test checks that the backend doesn't crash if the module has literal ;; This test checks that the backend doesn't crash if the module has literal
;; structs, i.e. structs whose type has no name. Typicaly clang generate such ;; structs, i.e. structs whose type has no name. Typicaly clang generate such
;; structs if the kernel contains OpenCL 2.0 blocks. The IR was produced with ;; structs if the kernel contains OpenCL 2.0 blocks. The IR was produced with
;; the following command: ;; the following command:
;; clang -cc1 -triple spir -cl-std=cl2.0 -O0 literal-struct.cl -emit-llvm -o test/literal-struct.ll ;; clang -cc1 -triple spir -cl-std=cl2.0 -O0 literal-struct.cl -emit-llvm -o test/literal-struct.ll
;; literal-struct.cl: ;; literal-struct.cl:
;; void foo() ;; void foo()
;; { ;; {
;; void (^myBlock)(void) = ^{}; ;; void (^myBlock)(void) = ^{};
;; myBlock(); ;; myBlock();
;; } ;; }
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -opaque-pointers=0 -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; CHECK: OpName %[[#StructType0:]] "struct.__opencl_block_literal_generic" ; CHECK: OpName %[[#StructType0:]] "struct.__opencl_block_literal_generic"
; CHECK: %[[#Int8:]] = OpTypeInt 8 0 ; CHECK: %[[#Int8:]] = OpTypeInt 8 0
; CHECK: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]] ; CHECK: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]]
; CHECK: %[[#Int:]] = OpTypeInt 32 0 ; CHECK: %[[#Int:]] = OpTypeInt 32 0
; CHECK: %[[#StructType0:]] = OpTypeStruct %[[#Int]] %[[#Int]] %[[#Int8Ptr]] ; CHECK: %[[#StructType0:]] = OpTypeStruct %[[#Int]] %[[#Int]] %[[#Int8Ptr]]
; CHECK: %[[#StructType:]] = OpTypeStruct %[[#Int]] %[[#Int]] %[[#Int8Ptr]] ; CHECK: %[[#StructType:]] = OpTypeStruct %[[#Int]] %[[#Int]] %[[#Int8Ptr]]
Show All 25 Lines

llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: OpDecorate %[[#NonConstMemset:]] LinkageAttributes "spirv.llvm_memset_p3i8_i32" ; CHECK-SPIRV: OpDecorate %[[#NonConstMemset:]] LinkageAttributes "spirv.llvm_memset_p3i8_i32"
; CHECK-SPIRV: %[[#Int32:]] = OpTypeInt 32 0 ; CHECK-SPIRV: %[[#Int32:]] = OpTypeInt 32 0
; CHECK-SPIRV: %[[#Int8:]] = OpTypeInt 8 0 ; CHECK-SPIRV: %[[#Int8:]] = OpTypeInt 8 0
; CHECK-SPIRV: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]] ; CHECK-SPIRV: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]]
; CHECK-SPIRV: %[[#Lenmemset21:]] = OpConstant %[[#]] 4 ; CHECK-SPIRV: %[[#Lenmemset21:]] = OpConstant %[[#]] 4
; CHECK-SPIRV: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Lenmemset21]] ; CHECK-SPIRV: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Lenmemset21]]
; CHECK-SPIRV: %[[#Int8PtrConst:]] = OpTypePointer UniformConstant %[[#Int8]] ; CHECK-SPIRV: %[[#Int8PtrConst:]] = OpTypePointer UniformConstant %[[#Int8]]
▲ Show 20 Lines • Show All 74 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/llvm-intrinsics/sqrt.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; CHECK: %[[#ExtInstSetId:]] = OpExtInstImport "OpenCL.std" ; CHECK: %[[#ExtInstSetId:]] = OpExtInstImport "OpenCL.std"
; CHECK: %[[#Float:]] = OpTypeFloat 32 ; CHECK: %[[#Float:]] = OpTypeFloat 32
; CHECK: %[[#Double:]] = OpTypeFloat 64 ; CHECK: %[[#Double:]] = OpTypeFloat 64
; CHECK: %[[#Double4:]] = OpTypeVector %[[#Double]] 4 ; CHECK: %[[#Double4:]] = OpTypeVector %[[#Double]] 4
; CHECK: %[[#FloatArg:]] = OpConstant %[[#Float]] ; CHECK: %[[#FloatArg:]] = OpConstant %[[#Float]]
; CHECK: %[[#DoubleArg:]] = OpConstant %[[#Double]] ; CHECK: %[[#DoubleArg:]] = OpConstant %[[#Double]]
; CHECK: %[[#Double4Arg:]] = OpConstantComposite %[[#Double4]] ; CHECK: %[[#Double4Arg:]] = OpConstantComposite %[[#Double4]]
Show All 21 Lines

llvm/test/CodeGen/SPIRV/opencl.queue_t.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: OpCapability DeviceEnqueue ; CHECK-SPIRV: OpCapability DeviceEnqueue
; CHECK-SPIRV: OpTypeQueue ; CHECK-SPIRV: OpTypeQueue
%opencl.queue_t = type opaque %opencl.queue_t = type opaque
define spir_func void @enqueue_simple_block(%opencl.queue_t* addrspace(3)* nocapture %q) { define spir_func void @enqueue_simple_block(%opencl.queue_t* addrspace(3)* nocapture %q) {
entry: entry:
ret void ret void
} }

llvm/test/CodeGen/SPIRV/opencl/basic/get_global_offset.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
; CHECK: OpEntryPoint Kernel %[[#test_func:]] "test" ; CHECK: OpEntryPoint Kernel %[[#test_func:]] "test"
; CHECK: OpName %[[#outOffsets:]] "outOffsets" ; CHECK: OpName %[[#outOffsets:]] "outOffsets"
; CHECK: OpName %[[#test_func]] "test" ; CHECK: OpName %[[#test_func]] "test"
; CHECK: OpName %[[#f2_decl:]] "BuiltInGlobalOffset" ; CHECK: OpName %[[#f2_decl:]] "BuiltInGlobalOffset"
; CHECK: OpDecorate %[[#f2_decl]] LinkageAttributes "BuiltInGlobalOffset" Import ; CHECK: OpDecorate %[[#f2_decl]] LinkageAttributes "BuiltInGlobalOffset" Import
; CHECK: %[[#int_ty:]] = OpTypeInt 32 0 ; CHECK: %[[#int_ty:]] = OpTypeInt 32 0
; CHECK: %[[#iptr_ty:]] = OpTypePointer CrossWorkgroup %[[#int_ty]] ; CHECK: %[[#iptr_ty:]] = OpTypePointer CrossWorkgroup %[[#int_ty]]
▲ Show 20 Lines • Show All 48 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_init.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
; CHECK: OpEntryPoint Kernel %[[#f1:]] "writer" ; CHECK: OpEntryPoint Kernel %[[#f1:]] "writer"
; CHECK: OpEntryPoint Kernel %[[#f2:]] "reader" ; CHECK: OpEntryPoint Kernel %[[#f2:]] "reader"
; CHECK-DAG: OpName %[[#a_var:]] "a_var" ; CHECK-DAG: OpName %[[#a_var:]] "a_var"
; CHECK-DAG: OpName %[[#p_var:]] "p_var" ; CHECK-DAG: OpName %[[#p_var:]] "p_var"
; CHECK-DAG: %[[#uchar:]] = OpTypeInt 8 0 ; CHECK-DAG: %[[#uchar:]] = OpTypeInt 8 0
; CHECK-DAG: %[[#pt1:]] = OpTypePointer CrossWorkgroup %[[#uchar]] ; CHECK-DAG: %[[#pt1:]] = OpTypePointer CrossWorkgroup %[[#uchar]]
; CHECK-DAG: %[[#arr2:]] = OpTypeArray %[[#uchar]] ; CHECK-DAG: %[[#arr2:]] = OpTypeArray %[[#uchar]]
▲ Show 20 Lines • Show All 100 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/opencl/basic/vstore_private.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
; CHECK: %[[#i16_ty:]] = OpTypeInt 16 0 ; CHECK: %[[#i16_ty:]] = OpTypeInt 16 0
; CHECK: %[[#v4xi16_ty:]] = OpTypeVector %[[#i16_ty]] 4 ; CHECK: %[[#v4xi16_ty:]] = OpTypeVector %[[#i16_ty]] 4
; CHECK: %[[#pv4xi16_ty:]] = OpTypePointer Function %[[#v4xi16_ty]] ; CHECK: %[[#pv4xi16_ty:]] = OpTypePointer Function %[[#v4xi16_ty]]
; CHECK: %[[#i16_const0:]] = OpConstant %[[#i16_ty]] 0 ; CHECK: %[[#i16_const0:]] = OpConstant %[[#i16_ty]] 0
; CHECK: %[[#i16_undef:]] = OpUndef %[[#i16_ty]] ; CHECK: %[[#i16_undef:]] = OpUndef %[[#i16_ty]]
; CHECK: %[[#comp_const:]] = OpConstantComposite %[[#v4xi16_ty]] %[[#i16_const0]] %[[#i16_const0]] %[[#i16_const0]] %[[#i16_undef]] ; CHECK: %[[#comp_const:]] = OpConstantComposite %[[#v4xi16_ty]] %[[#i16_const0]] %[[#i16_const0]] %[[#i16_const0]] %[[#i16_undef]]
▲ Show 20 Lines • Show All 83 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/opencl/image.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
;; FIXME: Write tests to ensure invalid usage of image are rejected, such as: ;; FIXME: Write tests to ensure invalid usage of image are rejected, such as:
;; - invalid AS (only global is allowed); ;; - invalid AS (only global is allowed);
;; - used in struct, union, array, pointer or return types; ;; - used in struct, union, array, pointer or return types;
;; - used with invalid CV-qualifiers (const or volatile in C99). ;; - used with invalid CV-qualifiers (const or volatile in C99).
;; FIXME: Write further tests to cover _array, _buffer, _depth, ... types. ;; FIXME: Write further tests to cover _array, _buffer, _depth, ... types.
%opencl.image1d_ro_t = type opaque ;; read_only image1d_t %opencl.image1d_ro_t = type opaque ;; read_only image1d_t
▲ Show 20 Lines • Show All 45 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/pstruct.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
%struct.ST = type { i32, i32, i32 } %struct.ST = type { i32, i32, i32 }
; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST" ; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST"
; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0 ; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0
; CHECK-SPIRV: %[[#intP:]] = OpTypePointer Function %[[#int]] ; CHECK-SPIRV: %[[#intP:]] = OpTypePointer Function %[[#int]]
; CHECK-SPIRV: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]] ; CHECK-SPIRV: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
; CHECK-SPIRV: %[[#structP:]] = OpTypePointer Function %[[#struct]] ; CHECK-SPIRV: %[[#structP:]] = OpTypePointer Function %[[#struct]]
▲ Show 20 Lines • Show All 112 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/read_image.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#IntTy:]] = OpTypeInt ; CHECK-SPIRV: %[[#IntTy:]] = OpTypeInt
; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]] ; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]]
; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat ; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat
; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]] ; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]]
; CHECK-SPIRV: OpImageRead %[[#IVecTy]] ; CHECK-SPIRV: OpImageRead %[[#IVecTy]]
; CHECK-SPIRV: OpImageRead %[[#FVecTy]] ; CHECK-SPIRV: OpImageRead %[[#FVecTy]]
▲ Show 20 Lines • Show All 41 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/spirv.Queue.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: OpCapability DeviceEnqueue ; CHECK-SPIRV: OpCapability DeviceEnqueue
; CHECK-SPIRV: OpTypeQueue ; CHECK-SPIRV: OpTypeQueue
%spirv.Queue = type opaque %spirv.Queue = type opaque
define spir_func void @enqueue_simple_block(%spirv.Queue* addrspace(3)* nocapture %q) { define spir_func void @enqueue_simple_block(%spirv.Queue* addrspace(3)* nocapture %q) {
entry: entry:
ret void ret void
} }

llvm/test/CodeGen/SPIRV/struct.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
%struct.ST = type { i32, i32, i32 } %struct.ST = type { i32, i32, i32 }
; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST" ; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST"
; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0 ; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0
; CHECK-SPIRV-DAG: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]] ; CHECK-SPIRV-DAG: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
; CHECK-SPIRV-DAG: %[[#structP:]] = OpTypePointer Function %[[#struct]] ; CHECK-SPIRV-DAG: %[[#structP:]] = OpTypePointer Function %[[#struct]]
; CHECK-SPIRV-DAG: %[[#intP:]] = OpTypePointer Function %[[#int]] ; CHECK-SPIRV-DAG: %[[#intP:]] = OpTypePointer Function %[[#int]]
Show All 37 Lines

llvm/test/CodeGen/SPIRV/transcoding/BuildNDRange_2.ll

Show All 12 Lines
;; const size_t lsize2[2] = {1, 1}; ;; const size_t lsize2[2] = {1, 1};
;; ndrange_2D(lsize2); ;; ndrange_2D(lsize2);
;; ;;
;; const size_t lsize3[3] = {1, 1, 1}; ;; const size_t lsize3[3] = {1, 1, 1};
;; ndrange_3D(lsize3); ;; ndrange_3D(lsize3);
;; } ;; }
;; bash$ $PATH_TO_GEN/bin/clang -cc1 -x cl -cl-std=CL2.0 -triple spir64-unknown-unknown -emit-llvm -include opencl-20.h BuildNDRange_2.cl -o BuildNDRange_2.ll ;; bash$ $PATH_TO_GEN/bin/clang -cc1 -x cl -cl-std=CL2.0 -triple spir64-unknown-unknown -emit-llvm -include opencl-20.h BuildNDRange_2.cl -o BuildNDRange_2.ll
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: %[[#LEN2_ID:]] = OpConstant %[[#]] 2 ; CHECK-SPIRV-DAG: %[[#LEN2_ID:]] = OpConstant %[[#]] 2
; CHECK-SPIRV-DAG: %[[#LEN3_ID:]] = OpConstant %[[#]] 3 ; CHECK-SPIRV-DAG: %[[#LEN3_ID:]] = OpConstant %[[#]] 3
; CHECK-SPIRV-DAG: %[[#ARRAY_T2:]] = OpTypeArray %[[#]] %[[#LEN2_ID]] ; CHECK-SPIRV-DAG: %[[#ARRAY_T2:]] = OpTypeArray %[[#]] %[[#LEN2_ID]]
; CHECK-SPIRV-DAG: %[[#ARRAY_T3:]] = OpTypeArray %[[#]] %[[#LEN3_ID]] ; CHECK-SPIRV-DAG: %[[#ARRAY_T3:]] = OpTypeArray %[[#]] %[[#LEN3_ID]]
; CHECK-SPIRV-LABEL: OpFunction ; CHECK-SPIRV-LABEL: OpFunction
; CHECK-SPIRV: %[[#LOAD2_ID:]] = OpLoad %[[#ARRAY_T2]] ; CHECK-SPIRV: %[[#LOAD2_ID:]] = OpLoad %[[#ARRAY_T2]]
▲ Show 20 Lines • Show All 49 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpConstantSampler.ll

;; Sources: ;; Sources:
;; ;;
;; void kernel foo(__read_only image2d_t src) { ;; void kernel foo(__read_only image2d_t src) {
;; sampler_t sampler1 = CLK_NORMALIZED_COORDS_TRUE | ;; sampler_t sampler1 = CLK_NORMALIZED_COORDS_TRUE |
;; CLK_ADDRESS_REPEAT | ;; CLK_ADDRESS_REPEAT |
;; CLK_FILTER_NEAREST; ;; CLK_FILTER_NEAREST;
;; sampler_t sampler2 = 0x00; ;; sampler_t sampler2 = 0x00;
;; ;;
;; read_imagef(src, sampler1, 0, 0); ;; read_imagef(src, sampler1, 0, 0);
;; read_imagef(src, sampler2, 0, 0); ;; read_imagef(src, sampler2, 0, 0);
;; } ;; }
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#SamplerID0:]] = OpConstantSampler %[[#]] Repeat 1 Nearest ; CHECK-SPIRV: %[[#SamplerID0:]] = OpConstantSampler %[[#]] Repeat 1 Nearest
; CHECK-SPIRV: %[[#SamplerID1:]] = OpConstantSampler %[[#]] None 0 Nearest ; CHECK-SPIRV: %[[#SamplerID1:]] = OpConstantSampler %[[#]] None 0 Nearest
; CHECK-SPIRV: %[[#]] = OpSampledImage %[[#]] %[[#]] %[[#SamplerID0]] ; CHECK-SPIRV: %[[#]] = OpSampledImage %[[#]] %[[#]] %[[#SamplerID0]]
; CHECK-SPIRV: %[[#]] = OpSampledImage %[[#]] %[[#]] %[[#SamplerID1]] ; CHECK-SPIRV: %[[#]] = OpSampledImage %[[#]] %[[#]] %[[#SamplerID1]]
%opencl.image2d_ro_t = type opaque %opencl.image2d_ro_t = type opaque
%opencl.sampler_t = type opaque %opencl.sampler_t = type opaque
Show All 14 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpImageQuerySize.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
;; Check conversion of get_image_width, get_image_height, get_image_depth, ;; Check conversion of get_image_width, get_image_height, get_image_depth,
;; get_image_array_size, and get_image_dim OCL built-ins. ;; get_image_array_size, and get_image_dim OCL built-ins.
;; In general the SPRI-V reader converts OpImageQuerySize into get_image_dim ;; In general the SPRI-V reader converts OpImageQuerySize into get_image_dim
;; and subsequent extract or shufflevector instructions. Unfortunately there is ;; and subsequent extract or shufflevector instructions. Unfortunately there is
;; no get_image_dim for 1D images and get_image_dim cannot replace get_image_array_size ;; no get_image_dim for 1D images and get_image_dim cannot replace get_image_array_size
; CHECK-SPIRV: %[[#ArrayTypeID:]] = OpTypeImage %[[#]] 1D 0 1 0 0 Unknown ReadOnly ; CHECK-SPIRV: %[[#ArrayTypeID:]] = OpTypeImage %[[#]] 1D 0 1 0 0 Unknown ReadOnly
▲ Show 20 Lines • Show All 117 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpImageReadMS.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#]] = OpImageRead %[[#]] %[[#]] %[[#]] Sample %[[#]] ; CHECK-SPIRV: %[[#]] = OpImageRead %[[#]] %[[#]] %[[#]] Sample %[[#]]
%opencl.image2d_msaa_ro_t = type opaque %opencl.image2d_msaa_ro_t = type opaque
define spir_kernel void @sample_test(%opencl.image2d_msaa_ro_t addrspace(1)* %source, i32 %sampler, <4 x float> addrspace(1)* nocapture %results) { define spir_kernel void @sample_test(%opencl.image2d_msaa_ro_t addrspace(1)* %source, i32 %sampler, <4 x float> addrspace(1)* nocapture %results) {
entry: entry:
%call = tail call spir_func i32 @_Z13get_global_idj(i32 0) %call = tail call spir_func i32 @_Z13get_global_idj(i32 0)
▲ Show 20 Lines • Show All 41 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpImageSampleExplicitLod.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: %[[#RetID:]] = OpImageSampleExplicitLod %[[#RetType:]] %[[#]] %[[#]] Lod %[[#]] ; CHECK-SPIRV-DAG: %[[#RetID:]] = OpImageSampleExplicitLod %[[#RetType:]] %[[#]] %[[#]] Lod %[[#]]
; CHECK-SPIRV-DAG: %[[#RetType]] = OpTypeVector %[[#]] 4 ; CHECK-SPIRV-DAG: %[[#RetType]] = OpTypeVector %[[#]] 4
; CHECK-SPIRV: %[[#]] = OpCompositeExtract %[[#]] %[[#RetID]] 0 ; CHECK-SPIRV: %[[#]] = OpCompositeExtract %[[#]] %[[#RetID]] 0
%opencl.image2d_depth_ro_t = type opaque %opencl.image2d_depth_ro_t = type opaque
define spir_kernel void @sample_kernel(%opencl.image2d_depth_ro_t addrspace(1)* %input, i32 %imageSampler, float addrspace(1)* %xOffsets, float addrspace(1)* %yOffsets, float addrspace(1)* %results) { define spir_kernel void @sample_kernel(%opencl.image2d_depth_ro_t addrspace(1)* %input, i32 %imageSampler, float addrspace(1)* %xOffsets, float addrspace(1)* %yOffsets, float addrspace(1)* %results) {
Show All 26 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpImageWrite.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#VOID_TY:]] = OpTypeVoid ; CHECK-SPIRV: %[[#VOID_TY:]] = OpTypeVoid
; CHECK-SPIRV: %[[#IMG2D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly ; CHECK-SPIRV: %[[#IMG2D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly
; CHECK-SPIRV: %[[#IMG2D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadWrite ; CHECK-SPIRV: %[[#IMG2D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadWrite
; CHECK-SPIRV: %[[#IMG2D_ARRAY_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 1 0 0 Unknown WriteOnly ; CHECK-SPIRV: %[[#IMG2D_ARRAY_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 1 0 0 Unknown WriteOnly
; CHECK-SPIRV: %[[#IMG2D_ARRAY_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 1 0 0 Unknown ReadWrite ; CHECK-SPIRV: %[[#IMG2D_ARRAY_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 1 0 0 Unknown ReadWrite
; CHECK-SPIRV: %[[#IMG1D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 0 0 0 Unknown WriteOnly ; CHECK-SPIRV: %[[#IMG1D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 0 0 0 Unknown WriteOnly
; CHECK-SPIRV: %[[#IMG1D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 0 0 0 Unknown ReadWrite ; CHECK-SPIRV: %[[#IMG1D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 0 0 0 Unknown ReadWrite
▲ Show 20 Lines • Show All 335 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_cmpxchg.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
;; This test checks that the backend is capable to correctly translate ;; This test checks that the backend is capable to correctly translate
;; atomic_cmpxchg OpenCL C 1.2 built-in function [1] into corresponding SPIR-V ;; atomic_cmpxchg OpenCL C 1.2 built-in function [1] into corresponding SPIR-V
;; instruction. ;; instruction.
;; __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) { ;; __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
;; atomic_cmpxchg(p, cmp, val); ;; atomic_cmpxchg(p, cmp, val);
;; ;;
▲ Show 20 Lines • Show All 43 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_legacy.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
;; This test checks that the backend is capable to correctly translate ;; This test checks that the backend is capable to correctly translate
;; legacy atomic OpenCL C 1.2 built-in functions [1] into corresponding SPIR-V ;; legacy atomic OpenCL C 1.2 built-in functions [1] into corresponding SPIR-V
;; instruction. ;; instruction.
;; __kernel void test_legacy_atomics(__global int *p, int val) { ;; __kernel void test_legacy_atomics(__global int *p, int val) {
;; atom_add(p, val); // from cl_khr_global_int32_base_atomics ;; atom_add(p, val); // from cl_khr_global_int32_base_atomics
;; atomic_add(p, val); // from OpenCL C 1.1 ;; atomic_add(p, val); // from OpenCL C 1.1
Show All 37 Lines

llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
;; constant sampler_t constSampl = CLK_FILTER_LINEAR; ;; constant sampler_t constSampl = CLK_FILTER_LINEAR;
;; ;;
;; __kernel ;; __kernel
;; void sample_kernel_float(image2d_t input, float2 coords, global float4 *results, sampler_t argSampl) { ;; void sample_kernel_float(image2d_t input, float2 coords, global float4 *results, sampler_t argSampl) {
;; *results = read_imagef(input, constSampl, coords); ;; *results = read_imagef(input, constSampl, coords);
;; *results = read_imagef(input, argSampl, coords); ;; *results = read_imagef(input, argSampl, coords);
;; *results = read_imagef(input, CLK_FILTER_NEAREST|CLK_ADDRESS_REPEAT, coords); ;; *results = read_imagef(input, CLK_FILTER_NEAREST|CLK_ADDRESS_REPEAT, coords);
▲ Show 20 Lines • Show All 80 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/block_w_struct_return.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4 ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4
;; TODO: We cannot check SPIR_V 1.1 and 1.4 simultaneously, implement additional ;; TODO: We cannot check SPIR_V 1.1 and 1.4 simultaneously, implement additional
;; run with CHECK-SPIRV1_1. ;; run with CHECK-SPIRV1_1.
;; kernel void block_ret_struct(__global int* res) ;; kernel void block_ret_struct(__global int* res)
;; { ;; {
;; struct A { ;; struct A {
;; int a; ;; int a;
▲ Show 20 Lines • Show All 98 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#IMAGE_TYPE:]] = OpTypeImage ; CHECK-SPIRV: %[[#IMAGE_TYPE:]] = OpTypeImage
; CHECK-SPIRV: %[[#IMAGE_ARG:]] = OpFunctionParameter %[[#IMAGE_TYPE]] ; CHECK-SPIRV: %[[#IMAGE_ARG:]] = OpFunctionParameter %[[#IMAGE_TYPE]]
; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#]] %[[#IMAGE_ARG]] ; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#]] %[[#IMAGE_ARG]]
; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#]] %[[#IMAGE_ARG]] ; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#]] %[[#IMAGE_ARG]]
%opencl.image2d_array_ro_t = type opaque %opencl.image2d_array_ro_t = type opaque
Show All 14 Lines

llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll

Show All 12 Lines
;; read_only image2d_array_t f1, ;; read_only image2d_array_t f1,
;; read_only image1d_buffer_t g1, ;; read_only image1d_buffer_t g1,
;; write_only image1d_t c2, ;; write_only image1d_t c2,
;; read_write image2d_t d3, ;; read_write image2d_t d3,
;; sampler_t s ;; sampler_t s
;; ) { ;; ) {
;; } ;; }
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: OpCapability Sampled1D ; CHECK-SPIRV-DAG: OpCapability Sampled1D
; CHECK-SPIRV-DAG: OpCapability SampledBuffer ; CHECK-SPIRV-DAG: OpCapability SampledBuffer
; CHECK-SPIRV-DAG: %[[#VOID:]] = OpTypeVoid ; CHECK-SPIRV-DAG: %[[#VOID:]] = OpTypeVoid
; CHECK-SPIRV-DAG: %[[#PIPE_RD:]] = OpTypePipe ReadOnly ; CHECK-SPIRV-DAG: %[[#PIPE_RD:]] = OpTypePipe ReadOnly
; CHECK-SPIRV-DAG: %[[#PIPE_WR:]] = OpTypePipe WriteOnly ; CHECK-SPIRV-DAG: %[[#PIPE_WR:]] = OpTypePipe WriteOnly
; CHECK-SPIRV-DAG: %[[#IMG1D_RD:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly ; CHECK-SPIRV-DAG: %[[#IMG1D_RD:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly
; CHECK-SPIRV-DAG: %[[#IMG2D_RD:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly ; CHECK-SPIRV-DAG: %[[#IMG2D_RD:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly
▲ Show 20 Lines • Show All 58 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/clk_event_t.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: OpTypeDeviceEvent ; CHECK-SPIRV: OpTypeDeviceEvent
; CHECK-SPIRV: OpFunction ; CHECK-SPIRV: OpFunction
; CHECK-SPIRV: OpCreateUserEvent ; CHECK-SPIRV: OpCreateUserEvent
; CHECK-SPIRV: OpIsValidEvent ; CHECK-SPIRV: OpIsValidEvent
; CHECK-SPIRV: OpRetainEvent ; CHECK-SPIRV: OpRetainEvent
; CHECK-SPIRV: OpSetUserEventStatus ; CHECK-SPIRV: OpSetUserEventStatus
; CHECK-SPIRV: OpCaptureEventProfilingInfo ; CHECK-SPIRV: OpCaptureEventProfilingInfo
Show All 38 Lines

llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer1:]] "__device_side_enqueue_block_invoke_kernel" ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer1:]] "__device_side_enqueue_block_invoke_kernel"
; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer2:]] "__device_side_enqueue_block_invoke_2_kernel" ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer2:]] "__device_side_enqueue_block_invoke_2_kernel"
; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer3:]] "__device_side_enqueue_block_invoke_3_kernel" ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer3:]] "__device_side_enqueue_block_invoke_3_kernel"
; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer4:]] "__device_side_enqueue_block_invoke_4_kernel" ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer4:]] "__device_side_enqueue_block_invoke_4_kernel"
; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer5:]] "__device_side_enqueue_block_invoke_5_kernel" ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer5:]] "__device_side_enqueue_block_invoke_5_kernel"
; CHECK-SPIRV: OpName %[[#BlockGlb1:]] "__block_literal_global" ; CHECK-SPIRV: OpName %[[#BlockGlb1:]] "__block_literal_global"
; CHECK-SPIRV: OpName %[[#BlockGlb2:]] "__block_literal_global.1" ; CHECK-SPIRV: OpName %[[#BlockGlb2:]] "__block_literal_global.1"
▲ Show 20 Lines • Show All 376 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/extract_insert_value.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
;; Check 'LLVM ==> SPIR-V' conversion of extractvalue/insertvalue. ;; Check 'LLVM ==> SPIR-V' conversion of extractvalue/insertvalue.
%struct.arr = type { [7 x float] } %struct.arr = type { [7 x float] }
%struct.st = type { %struct.inner } %struct.st = type { %struct.inner }
%struct.inner = type { float } %struct.inner = type { float }
; CHECK-SPIRV: %[[#float_ty:]] = OpTypeFloat 32 ; CHECK-SPIRV: %[[#float_ty:]] = OpTypeFloat 32
▲ Show 20 Lines • Show All 57 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/get_image_num_mip_levels.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
;; Types: ;; Types:
; CHECK-DAG: %[[#INT:]] = OpTypeInt 32 ; CHECK-DAG: %[[#INT:]] = OpTypeInt 32
; CHECK-DAG: %[[#VOID:]] = OpTypeVoid ; CHECK-DAG: %[[#VOID:]] = OpTypeVoid
; CHECK-DAG: %[[#IMAGE1D_T:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly ; CHECK-DAG: %[[#IMAGE1D_T:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly
; CHECK-DAG: %[[#IMAGE2D_T:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly ; CHECK-DAG: %[[#IMAGE2D_T:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly
; CHECK-DAG: %[[#IMAGE3D_T:]] = OpTypeImage %[[#VOID]] 3D 0 0 0 0 Unknown ReadOnly ; CHECK-DAG: %[[#IMAGE3D_T:]] = OpTypeImage %[[#VOID]] 3D 0 0 0 0 Unknown ReadOnly
; CHECK-DAG: %[[#IMAGE1D_ARRAY_T:]] = OpTypeImage %[[#VOID]] 1D 0 1 0 0 Unknown ReadOnly ; CHECK-DAG: %[[#IMAGE1D_ARRAY_T:]] = OpTypeImage %[[#VOID]] 1D 0 1 0 0 Unknown ReadOnly
▲ Show 20 Lines • Show All 73 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/global_block.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4 ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4
;; There are no blocks in SPIR-V. Therefore they are translated into regular ;; There are no blocks in SPIR-V. Therefore they are translated into regular
;; functions. An LLVM module which uses blocks, also contains some auxiliary ;; functions. An LLVM module which uses blocks, also contains some auxiliary
;; block-specific instructions, which are redundant in SPIR-V and should be ;; block-specific instructions, which are redundant in SPIR-V and should be
;; removed ;; removed
;; kernel void block_kernel(__global int* res) { ;; kernel void block_kernel(__global int* res) {
;; typedef int (^block_t)(int); ;; typedef int (^block_t)(int);
▲ Show 20 Lines • Show All 49 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/image_get_size_with_access_qualifiers.ll

; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: %[[#IntTyID:]] = OpTypeInt ; CHECK-SPIRV-DAG: %[[#IntTyID:]] = OpTypeInt
; CHECK-SPIRV-DAG: %[[#VoidTyID:]] = OpTypeVoid ; CHECK-SPIRV-DAG: %[[#VoidTyID:]] = OpTypeVoid
; CHECK-SPIRV-DAG: %[[#ImageTyID:]] = OpTypeImage %[[#VoidTyID]] 2D 0 1 0 0 Unknown ReadOnly ; CHECK-SPIRV-DAG: %[[#ImageTyID:]] = OpTypeImage %[[#VoidTyID]] 2D 0 1 0 0 Unknown ReadOnly
; CHECK-SPIRV-DAG: %[[#VectorTyID:]] = OpTypeVector %[[#IntTyID]] [[#]] ; CHECK-SPIRV-DAG: %[[#VectorTyID:]] = OpTypeVector %[[#IntTyID]] [[#]]
; CHECK-SPIRV: %[[#ImageArgID:]] = OpFunctionParameter %[[#ImageTyID]] ; CHECK-SPIRV: %[[#ImageArgID:]] = OpFunctionParameter %[[#ImageTyID]]
; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#VectorTyID]] %[[#ImageArgID]] ; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#VectorTyID]] %[[#ImageArgID]]
Show All 9 Lines

llvm/test/CodeGen/SPIRV/transcoding/image_with_access_qualifiers.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: OpCapability ImageReadWrite ; CHECK-SPIRV-DAG: OpCapability ImageReadWrite
; CHECK-SPIRV-DAG: OpCapability LiteralSampler ; CHECK-SPIRV-DAG: OpCapability LiteralSampler
; CHECK-SPIRV-DAG: %[[#TyVoid:]] = OpTypeVoid ; CHECK-SPIRV-DAG: %[[#TyVoid:]] = OpTypeVoid
; CHECK-SPIRV-DAG: %[[#TyImageID:]] = OpTypeImage %[[#TyVoid]] 1D 0 0 0 0 Unknown ReadWrite ; CHECK-SPIRV-DAG: %[[#TyImageID:]] = OpTypeImage %[[#TyVoid]] 1D 0 0 0 0 Unknown ReadWrite
; CHECK-SPIRV-DAG: %[[#TySampledImageID:]] = OpTypeSampledImage %[[#TyImageID]] ; CHECK-SPIRV-DAG: %[[#TySampledImageID:]] = OpTypeSampledImage %[[#TyImageID]]
; CHECK-SPIRV-DAG: %[[#ResID:]] = OpSampledImage %[[#TySampledImageID]] ; CHECK-SPIRV-DAG: %[[#ResID:]] = OpSampledImage %[[#TySampledImageID]]
Show All 14 Lines

llvm/test/CodeGen/SPIRV/transcoding/optional-core-features-multiple.ll

;; OpenCL C source ;; OpenCL C source
;; ----------------------------------------------- ;; -----------------------------------------------
;; double d = 1.0; ;; double d = 1.0;
;; kernel void test(read_only image2d_t img) {} ;; kernel void test(read_only image2d_t img) {}
;; ----------------------------------------------- ;; -----------------------------------------------
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
%opencl.image2d_t = type opaque %opencl.image2d_t = type opaque
@d = addrspace(1) global double 1.000000e+00, align 8 @d = addrspace(1) global double 1.000000e+00, align 8
define spir_kernel void @test(%opencl.image2d_t addrspace(1)* nocapture %img) { define spir_kernel void @test(%opencl.image2d_t addrspace(1)* nocapture %img) {
entry: entry:
ret void ret void
} }
; CHECK-SPIRV-DAG: OpCapability Float64 ; CHECK-SPIRV-DAG: OpCapability Float64
; CHECK-SPIRV-DAG: OpCapability ImageBasic ; CHECK-SPIRV-DAG: OpCapability ImageBasic

llvm/test/CodeGen/SPIRV/transcoding/spec_const.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-NOT: OpCapability Matrix ; CHECK-SPIRV-NOT: OpCapability Matrix
; CHECK-SPIRV-NOT: OpCapability Shader ; CHECK-SPIRV-NOT: OpCapability Shader
; CHECK-SPIRV: OpCapability Float16Buffer ; CHECK-SPIRV: OpCapability Float16Buffer
; CHECK-SPIRV-DAG: OpDecorate %[[#SC0:]] SpecId 0 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC0:]] SpecId 0
; CHECK-SPIRV-DAG: OpDecorate %[[#SC1:]] SpecId 1 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC1:]] SpecId 1
; CHECK-SPIRV-DAG: OpDecorate %[[#SC2:]] SpecId 2 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC2:]] SpecId 2
▲ Show 20 Lines • Show All 52 Lines • Show Last 20 Lines

llvm/test/CodeGen/SPIRV/transcoding/spirv-private-array-initialization.ll

; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; ;
; CHECK-SPIRV-DAG: %[[#i32:]] = OpTypeInt 32 0 ; CHECK-SPIRV-DAG: %[[#i32:]] = OpTypeInt 32 0
; CHECK-SPIRV-DAG: %[[#i8:]] = OpTypeInt 8 0 ; CHECK-SPIRV-DAG: %[[#i8:]] = OpTypeInt 8 0
; CHECK-SPIRV-DAG: %[[#one:]] = OpConstant %[[#i32]] 1 ; CHECK-SPIRV-DAG: %[[#one:]] = OpConstant %[[#i32]] 1
; CHECK-SPIRV-DAG: %[[#two:]] = OpConstant %[[#i32]] 2 ; CHECK-SPIRV-DAG: %[[#two:]] = OpConstant %[[#i32]] 2
; CHECK-SPIRV-DAG: %[[#three:]] = OpConstant %[[#i32]] 3 ; CHECK-SPIRV-DAG: %[[#three:]] = OpConstant %[[#i32]] 3
; CHECK-SPIRV-DAG: %[[#i32x3:]] = OpTypeArray %[[#i32]] %[[#three]] ; CHECK-SPIRV-DAG: %[[#i32x3:]] = OpTypeArray %[[#i32]] %[[#three]]
; CHECK-SPIRV-DAG: %[[#i32x3_ptr:]] = OpTypePointer Function %[[#i32x3]] ; CHECK-SPIRV-DAG: %[[#i32x3_ptr:]] = OpTypePointer Function %[[#i32x3]]
Show All 34 Lines

llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll

;; Test SPIR-V opaque types ;; Test SPIR-V opaque types
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV-DAG: OpCapability Float16 ; CHECK-SPIRV-DAG: OpCapability Float16
; CHECK-SPIRV-DAG: OpCapability ImageReadWrite ; CHECK-SPIRV-DAG: OpCapability ImageReadWrite
; CHECK-SPIRV-DAG: OpCapability Pipes ; CHECK-SPIRV-DAG: OpCapability Pipes
; CHECK-SPIRV-DAG: OpCapability DeviceEnqueue ; CHECK-SPIRV-DAG: OpCapability DeviceEnqueue
; CHECK-SPIRV-DAG: %[[#VOID:]] = OpTypeVoid ; CHECK-SPIRV-DAG: %[[#VOID:]] = OpTypeVoid
; CHECK-SPIRV-DAG: %[[#INT:]] = OpTypeInt 32 0 ; CHECK-SPIRV-DAG: %[[#INT:]] = OpTypeInt 32 0
▲ Show 20 Lines • Show All 91 Lines • Show Last 20 Lines