Skip to content

Commit

Permalink
Fix remaining XFAILs (google#1306)
Browse files Browse the repository at this point in the history 
ref google#1292

* Update tests to remove XFAILs
  • Loading branch information
@alan-baker
alan-baker authored Feb 9, 2024
1 parent 1907bf6 commit e7a0c9b
Show file tree
Hide file tree
Showing 5 changed files with 79 additions and 98 deletions.
5 changes: 1 addition & 4 deletions test/DirectResourceAccess/partial_access_chain_global.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,17 +3,14 @@
// RUN: FileCheck %s < %t2.spvasm
// RUN: spirv-val --target-env vulkan1.0 %t.spv

// TODO(#1292)
// XFAIL: *

// Kernel |bar| does a non-trivial access chain before calling the helper.

__attribute__((noinline))
void apple(global int *A, global int *B, int n) { A[n] = B[n + 2]; }

kernel void foo(global int *A, global int *B, int n) { apple(A, B, n); }

kernel void bar(global int *A, global int *B, int n) { apple(A + 1, B, n); }
kernel void bar(global int *A, global int *B, int n) { apple(A + n, B, n); }
// CHECK: OpEntryPoint GLCompute [[_33:%[0-9a-zA-Z_]+]] "foo"
// CHECK: OpEntryPoint GLCompute [[_40:%[0-9a-zA-Z_]+]] "bar"
// CHECK-DAG: OpDecorate [[_21:%[0-9a-zA-Z_]+]] Binding 1
Expand Down
40 changes: 0 additions & 40 deletions test/LongVectorLowering/bitselect_float8.cl
View file

This file was deleted.

67 changes: 67 additions & 0 deletions test/LongVectorLowering/bitselect_float8.ll
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,67 @@
; RUN: clspv-opt %s -o %t.ll --long-vector --passes=long-vector-lowering
; RUN: FileCheck %s < %t.ll

; CHECK-COUNT-8: extractvalue [8 x float]
; CHECK-COUNT-8: bitcast float %{{.*}} to i32
; CHECK-COUNT-8: xor i32 %{{.*}}, -1
; CHECK-COUNT-8: and i32
; CHECK-COUNT-8: and i32
; CHECK-COUNT-8: or i32
; CHECK-COUNT-8: bitcast i32 %{{.*}} to float

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-unknown-unknown"

@__spirv_WorkgroupSize = addrspace(8) global <3 x i32> zeroinitializer

define spir_kernel void @test(ptr addrspace(1) align 4 %a, ptr addrspace(1) align 4 %b, ptr addrspace(1) align 4 %c, ptr addrspace(1) align 4 %d, { i32 } %podargs) !clspv.pod_args_impl !13 !kernel_arg_map !14 {
entry:
%x = extractvalue { i32 } %podargs, 0
%arrayidx.i = getelementptr inbounds <8 x float>, ptr addrspace(1) %a, i32 %x
%0 = load <8 x float>, ptr addrspace(1) %arrayidx.i, align 4
%arrayidx1.i = getelementptr inbounds <8 x float>, ptr addrspace(1) %b, i32 %x
%1 = load <8 x float>, ptr addrspace(1) %arrayidx1.i, align 4
%arrayidx2.i = getelementptr inbounds <8 x float>, ptr addrspace(1) %c, i32 %x
%2 = load <8 x float>, ptr addrspace(1) %arrayidx2.i, align 4
%3 = bitcast <8 x float> %2 to <8 x i32>
%4 = bitcast <8 x float> %0 to <8 x i32>
%5 = bitcast <8 x float> %1 to <8 x i32>
%6 = xor <8 x i32> %3, <i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1>
%7 = and <8 x i32> %6, %4
%8 = and <8 x i32> %3, %5
%9 = or <8 x i32> %7, %8
%10 = bitcast <8 x i32> %9 to <8 x float>
%arrayidx3.i = getelementptr inbounds <8 x float>, ptr addrspace(1) %d, i32 %x
store <8 x float> %10, ptr addrspace(1) %arrayidx3.i, align 4
ret void
}

attributes #0 = { convergent norecurse nounwind "no-builtins" "no-trapping-math"="true" "stack-protector-buffer-size"="0" "stackrealign" "uniform-work-group-size"="true" }

!llvm.module.flags = !{!0, !1, !2}
!opencl.ocl.version = !{!3}
!opencl.spir.version = !{!3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3, !3}
!llvm.ident = !{!4, !5, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !5, !5, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6}
!_Z28clspv.entry_point_attributes = !{!7}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 7, !"direct-access-external-data", i32 0}
!2 = !{i32 7, !"frame-pointer", i32 2}
!3 = !{i32 1, i32 2}
!4 = !{!"clang version 19.0.0git (https://github.com/llvm/llvm-project d5a3de4aeef4f4f1c52692533ddb9fdf45aef9d3)"}
!5 = !{!"clang version 17.0.0 (https://github.com/llvm/llvm-project 1e6fc9626c0f49ce952a67aef47e86253d13f74a)"}
!6 = !{!"clang version 17.0.0 (https://github.com/llvm/llvm-project ab674234c440ed27302f58eeccc612c83b32c43f)"}
!7 = !{!"test", !" kernel"}
!8 = !{i32 1, i32 1, i32 1, i32 1, i32 0}
!9 = !{!"none", !"none", !"none", !"none", !"none"}
!10 = !{!"float*", !"float*", !"float*", !"float*", !"int"}
!11 = !{!"", !"", !"", !"", !""}
!12 = !{!"a", !"b", !"c", !"d", !"x"}
!13 = !{i32 2}
!14 = !{!15, !16, !17, !18, !19}
!15 = !{!"a", i32 0, i32 0, i32 0, i32 0, !"buffer"}
!16 = !{!"b", i32 1, i32 1, i32 0, i32 0, !"buffer"}
!17 = !{!"c", i32 2, i32 2, i32 0, i32 0, !"buffer"}
!18 = !{!"d", i32 3, i32 3, i32 0, i32 0, !"buffer"}
!19 = !{!"x", i32 4, i32 4, i32 0, i32 4, !"pod_pushconstant"}

7 changes: 2 additions & 5 deletions test/PointerAccessChains/pointer_array_stride_16.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,9 +3,6 @@
// RUN: FileCheck %s < %t.spvasm
// RUN: spirv-val --target-env vulkan1.0 %t.spv

// TODO(#1292)
// XFAIL: *

struct A {
float4 x;
};
Expand All @@ -19,8 +16,8 @@ static float4 bar(global struct A* in, int n) {
return in[n].x;
}

kernel void foo(global float* out, global struct B* in, int n) {
*out = bar(&in->a[1], n)[0];
kernel void foo(global float* out, global struct B* in, int n, int x) {
*out = bar(&in[x].a[1], n)[0];
}

// CHECK-DAG: OpDecorate [[array_struct_A:%[a-zA-Z0-9_]+]] ArrayStride 16
Expand Down
58 changes: 9 additions & 49 deletions test/RewritePackedStructs/packed_struct.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,9 +3,6 @@
// RUN: FileCheck %s < %t2.spvasm
// RUN: spirv-val --target-env vulkan1.0 %t.spv

// TODO(#1292)
// XFAIL: *

struct S1{
int x;
char y;
Expand Down Expand Up @@ -53,49 +50,12 @@ __kernel void test3(__global struct S3* a, __global struct S4* b) {
b[0].c = b[0].a + b[0].b + b[0].y;
}

// CHECK-DAG: [[uint:%[^ ]+]] = OpTypeInt 32 0
// CHECK-DAG: [[uint5:%[^ ]+]] = OpConstant [[uint]] 5
// CHECK-DAG: [[uint7:%[^ ]+]] = OpConstant [[uint]] 7
// CHECK-DAG: [[uint8:%[^ ]+]] = OpConstant [[uint]] 8
// CHECK-DAG: [[uchar:%[^ ]+]] = OpTypeInt 8 0

// CHECK-DAG: [[arr_uchar5:%[^ ]+]] = OpTypeArray [[uchar]] [[uint5]]
// CHECK-DAG: [[struct_arr_uchar5:%[^ ]+]] = OpTypeStruct [[arr_uchar5]]
// CHECK-DAG: [[arr_struct_arr_uchar5:%[^ ]+]] = OpTypeRuntimeArray [[struct_arr_uchar5]]
// CHECK-DAG: [[S1:%[^ ]+]] = OpTypeStruct [[arr_struct_arr_uchar5]]
// CHECK-DAG: [[S1_ptr:%[^ ]+]] = OpTypePointer StorageBuffer [[S1]]

// CHECK-DAG: [[arr_uchar8:%[^ ]+]] = OpTypeArray [[uchar]] [[uint8]]
// CHECK-DAG: [[arr_arr_uchar8:%[^ ]+]] = OpTypeRuntimeArray [[arr_uchar8]]
// CHECK-DAG: [[S2_4:%[^ ]+]] = OpTypeStruct [[arr_arr_uchar8]]
// CHECK-DAG: [[S2_4_ptr:%[^ ]+]] = OpTypePointer StorageBuffer [[S2_4]]

// CHECK-DAG: [[arr_uchar7:%[^ ]+]] = OpTypeArray [[uchar]] [[uint7]]
// CHECK-DAG: [[struct_arr_uchar7:%[^ ]+]] = OpTypeStruct [[arr_uchar7]]
// CHECK-DAG: [[arr_struct_arr_uchar7:%[^ ]+]] = OpTypeRuntimeArray [[struct_arr_uchar7]]
// CHECK-DAG: [[S3:%[^ ]+]] = OpTypeStruct [[arr_struct_arr_uchar7]]
// CHECK-DAG: [[S3_ptr:%[^ ]+]] = OpTypePointer StorageBuffer [[S3]]

// CHECK-DAG: OpDecorate [[arr_uchar5]] ArrayStride 1
// CHECK-DAG: OpDecorate [[arr_uchar7]] ArrayStride 1
// CHECK-DAG: OpDecorate [[arr_uchar8]] ArrayStride 1

// CHECK-DAG: OpMemberDecorate [[struct_arr_uchar5]] 0 Offset 0
// CHECK-DAG: OpDecorate [[arr_struct_arr_uchar5]] ArrayStride 5
// CHECK-DAG: OpDecorate [[arr_arr_uchar8]] ArrayStride 8
// CHECK-DAG: OpMemberDecorate [[struct_arr_uchar7]] 0 Offset 0
// CHECK-DAG: OpDecorate [[arr_struct_arr_uchar7]] ArrayStride 7

// CHECK-DAG: [[a_S1:%[^ ]+]] = OpVariable [[S1_ptr]] StorageBuffer
// CHECK-DAG: [[b_S2_4:%[^ ]+]] = OpVariable [[S2_4_ptr]] StorageBuffer
// CHECK-DAG: [[c_S3:%[^ ]+]] = OpVariable [[S3_ptr]] StorageBuffer
// CHECK-DAG: [[a_S3:%[^ ]+]] = OpVariable [[S3_ptr]] StorageBuffer

// CHECK-DAG: OpDecorate [[a_S1]] DescriptorSet 0
// CHECK-DAG: OpDecorate [[a_S1]] Binding 0
// CHECK-DAG: OpDecorate [[b_S2_4]] DescriptorSet 0
// CHECK-DAG: OpDecorate [[b_S2_4]] Binding 1
// CHECK-DAG: OpDecorate [[c_S3]] DescriptorSet 0
// CHECK-DAG: OpDecorate [[c_S3]] Binding 2
// CHECK-DAG: OpDecorate [[a_S3]] DescriptorSet 0
// CHECK-DAG: OpDecorate [[a_S3]] Binding 0
// CHECK-DAG: [[uint:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0
// CHECK-DAG: [[uchar:%[a-zA-Z0-9_]+]] = OpTypeInt 8 0
// CHECK-DAG: [[char_array:%[a-zA-Z0-9_]+]] = OpTypeRuntimeArray [[uchar]]
// CHECK-DAG: [[block:%[a-zA-Z0-9_]+]] = OpTypeStruct [[char_array]]
// CHECK-DAG: [[block_ptr:%[a-zA-Z0-9_]+]] = OpTypePointer StorageBuffer [[block]]
// CHECK-DAG: OpDecorate [[char_array]] ArrayStride 1
// CHECK-DAG: OpVariable [[block_ptr]] StorageBuffer
// CHECK-DAG: OpVariable [[block_ptr]] StorageBuffer
// CHECK-DAG: OpVariable [[block_ptr]] StorageBuffer

0 comments on commit e7a0c9b

Please sign in to comment.