Skip to content

[NFC] Add SPIR-V backend flow to transcoding tests #3444

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
MrSidims wants to merge 1 commit into
base: main
Choose a base branch
from
Draft

[NFC] Add SPIR-V backend flow to transcoding tests #3444

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions test/transcoding/AllowIntrinsics.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; Check that passing a prefix list into the option works/fails as expected.
;
; Positive cases:
Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,13 +6,21 @@
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// TODO: SPIR-V backend test disabled due to failures
// RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=CL2.0 -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

// RUN: llvm-spirv %t.bc -spirv-text -o %t.txt --spirv-ext=+SPV_KHR_untyped_pointers
// RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV
// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
// RUN: spirv-val %t.spv
// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// TODO: SPIR-V backend test disabled due to failures
// RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=CL2.0 -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

#define DEFINE_KERNEL(TYPE) \
__kernel void testAtomicCompareExchangeExplicit_cl20_##TYPE( \
volatile global atomic_##TYPE* object, \
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/BuiltinPrintf.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,10 @@
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// Test SPIR-V backend:
// RUN: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

// CHECK-SPIRV: ExtInst [[#]] [[#]] [[#]] printf [[#]]
// CHECK-LLVM: call spir_func i32 (ptr addrspace(2), ...) @printf(ptr addrspace(2) {{.*}})

Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/CommonAddrspaceGlobalVar.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-LLVM-NOT: alloca
; CHECK-LLVM: @CAG = common addrspace(1) global i32 0, align 4
; CHECK-LLVM-NOT: alloca
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/ConvertPtr.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,10 @@
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// Test SPIR-V backend:
// RUN: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

kernel void testConvertPtrToU(global int *a, global unsigned long *res) {
res[0] = (unsigned long)&a[0];
}
Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/CreatePipeFromPipeStorage.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,12 +5,20 @@
; RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=SPV-IR -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-SPV-IR

; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
; RUN: llvm-spirv %t.bc -spirv-text -o %t.txt --spirv-ext=+SPV_KHR_untyped_pointers
; RUN: FileCheck < %t.txt %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-UNTYPED-PTR
; RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=SPV-IR -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-SPV-IR


; CHECK-SPV-IR: %"[[CL_PIPE_STORAGE_NAME:[^"]+]]" = type { ptr addrspace(1) }
; CHECK-SPV-IR: %"[[CL_READ_PIPE_NAME:[^"]+read>]]" = type { target("spirv.Pipe", 0) }
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/DecorationAlignment.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
target triple = "spir64-unknown-unknown"

Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/DecorationMaxByteOffset.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,13 +4,21 @@
; RUN: llvm-spirv %t.bc -o %t.spv
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
; RUN: llvm-spirv %t.bc -spirv-text --spirv-max-version=1.0 -o - | FileCheck %s --check-prefix=CHECK-SPIRV_1_0

; RUN: llvm-spirv %t.bc -spirv-text -o %t.txt --spirv-ext=+SPV_KHR_untyped_pointers
; RUN: FileCheck < %t.txt %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-UNTYPED-PTR
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
; RUN: llvm-spirv %t.bc -spirv-text --spirv-max-version=1.0 --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s --check-prefix=CHECK-SPIRV_1_0


Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/DivRem.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,10 @@
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// Test SPIR-V backend:
// RUN: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

// CHECK-SPIRV-DAG: Name [[#__clang_ocl_kern_imp_testSDiv:]] "__clang_ocl_kern_imp_testSDiv"
// CHECK-SPIRV-DAG: Name [[#__clang_ocl_kern_imp_testUDiv:]] "__clang_ocl_kern_imp_testUDiv"
// CHECK-SPIRV-DAG: Name [[#__clang_ocl_kern_imp_testFDiv:]] "__clang_ocl_kern_imp_testFDiv"
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/ExecutionMode_SPIR_to_SPIRV.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; LLVM => SPIRV checks
; check for magic number followed by version 1.0
; CHECK-SPIRV: 119734787 65536
Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/OpAllAny.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,9 +5,17 @@
; RUN: spirv-val %t.spv
; RUN: llvm-spirv -r %t.spv -o %t.rev.spv.bc --spirv-target-env=SPV-IR
; RUN: llvm-dis < %t.rev.spv.bc | FileCheck %s --check-prefix=CHECK-SPV-IR

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=SPV-IR -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-SPV-IR
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; This test checks SYCL relational builtin any and all with vector input types.

; CHECK-SPIRV: 2 TypeBool [[BoolTypeID:[0-9]+]]
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpBitReverse_i32.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,10 @@
; RUN: spirv-val %t.spv
; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: 4 TypeInt [[int:[0-9]+]] 32
; CHECK-SPIRV: 4 BitReverse [[int]]

Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpBitReverse_v2i16.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,10 @@
; RUN: spirv-val %t.spv
; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: 4 TypeInt [[short:[0-9]+]] 16
; CHECK-SPIRV: 4 TypeVector [[short2:[0-9]+]] [[short]] 2
; CHECK-SPIRV: 4 BitReverse [[short2]]
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpConstantBool.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-LLVM-LABEL: @f
; CHECK-LLVM: ret i1 true
; CHECK-LLVM-LABEL: @f2
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpConstantSampler.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: ConstantSampler {{[0-9]+}} [[SamplerID0:[0-9]+]] 3 1 0
; CHECK-SPIRV: ConstantSampler {{[0-9]+}} [[SamplerID1:[0-9]+]] 0 0 0
; CHECK-SPIRV: SampledImage {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} [[SamplerID0]]
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpDot.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; The OpDot operands must be vectors; check that translating dot with
; scalar arguments does not result in OpDot.
; CHECK-SPIRV-LABEL: 5 Function
Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/OpGenericPtrMemSemantics.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,9 +5,17 @@
; RUN: spirv-val %t.spv
; RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=SPV-IR -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-SPV-IR
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: 4 GenericPtrMemSemantics {{[0-9]+}} [[ResID:[0-9]+]] {{[0-9]+}}
; CHECK-SPIRV-NEXT: 5 ShiftRightLogical {{[0-9]+}} {{[0-9]+}} [[ResID]] {{[0-9]+}}

Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpGroupAllAny.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-LLVM: call spir_func i32 @_Z14work_group_alli(
; CHECK-LLVM: call spir_func i32 @_Z14work_group_anyi(

Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/OpGroupAsyncCopy.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,9 +5,17 @@
; RUN: spirv-val %t.spv
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
; RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=SPV-IR -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-SPV-IR

; CHECK-LLVM: call spir_func target("spirv.Event") @_Z29async_work_group_strided_copyPU3AS1Dv2_hPU3AS3KS_jj9ocl_event(
; CHECK-LLVM: call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(
; CHECK-LLVM: declare spir_func target("spirv.Event") @_Z29async_work_group_strided_copyPU3AS1Dv2_hPU3AS3KS_jj9ocl_event(ptr addrspace(1), ptr addrspace(3), i32, i32, target("spirv.Event"))
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpImageReadMS.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-LLVM: call spir_func <4 x float> @_Z11read_imagef19ocl_image2d_msaa_roDv2_ii(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0)

; CHECK-SPIRV: 7 ImageRead {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} 64 {{[0-9]+}}
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpImageSampleExplicitLod.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-LLVM: call spir_func float @_Z11read_imagef20ocl_image2d_depth_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)

; CHECK-SPIRV-DAG: 7 ImageSampleExplicitLod [[RetType:[0-9]+]] [[RetID:[0-9]+]] {{[0-9]+}} {{[0-9]+}} 2 {{[0-9]+}}
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpImageSampleExplicitLod_arg.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// TODO: SPIR-V backend test disabled due to failures
// RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

void __kernel sample_kernel_read( __global float4 *results,
read_only image2d_t image,
sampler_t imageSampler,
Expand Down
8 changes: 8 additions & 0 deletions test/transcoding/OpImageWrite.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,9 +4,17 @@
// RUN: llvm-spirv %t.bc -o %t.spv
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

// TODO: SPIR-V backend test disabled due to failures
// RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc --spirv-target-env=SPV-IR
// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR

// TODO: SPIR-V backend test disabled due to failures
// RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
// RUNx: llvm-spirv -r %t.llc.spv --spirv-target-env=SPV-IR -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-SPV-IR

// CHECK-SPIRV: TypeVoid [[VOID_TY:[0-9]+]]
// CHECK-SPIRV: TypeImage [[IMG2D_WO_TY:[0-9]+]] [[VOID_TY]] 1 0 0 0 0 0 1
// CHECK-SPIRV: TypeImage [[IMG2D_RW_TY:[0-9]+]] [[VOID_TY]] 1 0 0 0 0 0 2
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpLine.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,10 @@
; RUN: llvm-spirv -to-text %t.spv -o -| FileCheck %s --check-prefix=CHECK-SPIRV
; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; XFAIL: *
; This is requires debug metadata update.
; We also should remove checks for LLVM IR stuff like labels, names, etc.
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpMin.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: 5 ExtInstImport [[SetInstID:[0-9]+]] "OpenCL.std"
; CHECK-SPIRV: 4 TypeInt [[IntTypeID:[0-9]+]] 32 {{[0-9]+}}
; CHECK-SPIRV: 4 TypeVector [[Int2TypeID:[0-9]+]] [[IntTypeID]] 2
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpPhi_ArgumentsPlaceholders.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
target triple = "spir"

Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpSwitch32.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.bc
; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: 7 Switch {{[0-9]+}} {{[0-9]+}} 0 {{[0-9]+}} 1 {{[0-9]+}}

; ModuleID = 'main'
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpSwitch64.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.bc
; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM

; TODO: SPIR-V backend test disabled due to failures
; RUNx: llc -O0 -mtriple=spirv64-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUNx: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: 12 Switch {{[0-9]+}} {{[0-9]+}} 0 0 {{[0-9]+}} 1 0 {{[0-9]+}} 1 5 {{[0-9]+}}

; ModuleID = 'main'
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpSwitchEmpty.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,10 @@
; RUN: spirv-val %t.spv
; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-SPIRV: FunctionParameter {{[0-9]+}} [[X:[0-9]+]]
; CHECK-SPIRV: Switch [[X]] [[DEFAULT:[0-9]+]] {{$}}
; CHECK-SPIRV: Label [[DEFAULT]]
Expand Down
4 changes: 4 additions & 0 deletions test/transcoding/OpVariable_Initializer.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM

; Test SPIR-V backend:
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %t.bc -o %t.llc.spv -filetype=obj
; RUN: llvm-spirv -r %t.llc.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM

; CHECK-LLVM: @test_atomic_fn.L = internal addrspace(3) global [64 x i32] poison, align 4

; CHECK-SPIRV-NOT: undef
Expand Down
Loading
Loading