From e7a0c9b97e7c38d3b6a7b744b0d39743d318dd2b Mon Sep 17 00:00:00 2001 From: alan-baker Date: Fri, 9 Feb 2024 02:35:43 -0500 Subject: [PATCH] Fix remaining XFAILs (#1306) ref #1292 * Update tests to remove XFAILs --- .../partial_access_chain_global.cl | 5 +- test/LongVectorLowering/bitselect_float8.cl | 40 ----------- test/LongVectorLowering/bitselect_float8.ll | 67 +++++++++++++++++++ .../pointer_array_stride_16.cl | 7 +- test/RewritePackedStructs/packed_struct.cl | 58 +++------------- 5 files changed, 79 insertions(+), 98 deletions(-) delete mode 100644 test/LongVectorLowering/bitselect_float8.cl create mode 100644 test/LongVectorLowering/bitselect_float8.ll diff --git a/test/DirectResourceAccess/partial_access_chain_global.cl b/test/DirectResourceAccess/partial_access_chain_global.cl index c24002b0c..e87fa2460 100644 --- a/test/DirectResourceAccess/partial_access_chain_global.cl +++ b/test/DirectResourceAccess/partial_access_chain_global.cl @@ -3,9 +3,6 @@ // 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)) @@ -13,7 +10,7 @@ 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 diff --git a/test/LongVectorLowering/bitselect_float8.cl b/test/LongVectorLowering/bitselect_float8.cl deleted file mode 100644 index d2be365e8..000000000 --- a/test/LongVectorLowering/bitselect_float8.cl +++ /dev/null @@ -1,40 +0,0 @@ -// RUN: clspv %target %s --long-vector -o %t.spv -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// TODO(#1292) -// XFAIL: * - -__kernel void test_bitselect(__global float8 *A, __global float8 *B, - __global float8 *C, __global float8 *destValue) { - *destValue = bitselect(*A, *B, *C); -} - -// CHECK-DAG: %[[uint:[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: %[[float:[0-9a-zA-Z_]+]] = OpTypeFloat 32 -// CHECK-DAG: %[[uint_8:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 8 -// CHECK-DAG: %[[float8:[0-9a-zA-Z_]+]] = OpTypeArray %[[float]] %[[uint_8]] -// CHECK-DAG: %[[array_float8:[0-9a-zA-Z_]+]] = OpTypeRuntimeArray %[[float8]] -// CHECK-DAG: %[[struct_float8:[0-9a-zA-Z_]+]] = OpTypeStruct %[[array_float8]] -// CHECK-DAG: %[[ptr_float8:[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer %[[struct_float8]] -// CHECK-DAG: %[[ptr_float:[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer %[[float]] -// CHECK-DAG: %[[uint0:[0-9a-zA-Z_]+]] = OpConstant %[[uint]] 0 - -// CHECK: %[[A:[0-9a-zA-Z_]+]] = OpVariable %[[ptr_float8]] StorageBuffer -// CHECK: %[[B:[0-9a-zA-Z_]+]] = OpVariable %[[ptr_float8]] StorageBuffer -// CHECK: %[[C:[0-9a-zA-Z_]+]] = OpVariable %[[ptr_float8]] StorageBuffer -// CHECK: %[[destValue:[0-9a-zA-Z_]+]] = OpVariable %[[ptr_float8]] StorageBuffer - -// CHECK: %[[GEPA:[0-9a-zA-Z_]+]] = OpAccessChain %[[ptr_float]] %[[A]] %[[uint0]] %[[uint0]] %[[uint0]] -// CHECK: %[[A0:[0-9a-zA-Z_]+]] = OpLoad %[[float]] %[[GEPA]] -// CHECK: %[[A0_uint:[0-9a-zA-Z_]+]] = OpBitcast %[[uint]] %[[A0]] - -// CHECK: %[[GEPB:[0-9a-zA-Z_]+]] = OpAccessChain %[[ptr_float]] %[[B]] %[[uint0]] %[[uint0]] %[[uint0]] -// CHECK: %[[B0:[0-9a-zA-Z_]+]] = OpLoad %[[float]] %[[GEPB]] -// CHECK: %[[B0_uint:[0-9a-zA-Z_]+]] = OpBitcast %[[uint]] %[[B0]] - -// CHECK: %[[GEPC:[0-9a-zA-Z_]+]] = OpAccessChain %[[ptr_float]] %[[C]] %[[uint0]] %[[uint0]] %[[uint0]] -// CHECK: %[[C0:[0-9a-zA-Z_]+]] = OpLoad %[[float]] %[[GEPC]] -// CHECK: %[[C0_uint:[0-9a-zA-Z_]+]] = OpBitcast %[[uint]] %[[C0]] - diff --git a/test/LongVectorLowering/bitselect_float8.ll b/test/LongVectorLowering/bitselect_float8.ll new file mode 100644 index 000000000..9400bcd5e --- /dev/null +++ b/test/LongVectorLowering/bitselect_float8.ll @@ -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, + %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"} + diff --git a/test/PointerAccessChains/pointer_array_stride_16.cl b/test/PointerAccessChains/pointer_array_stride_16.cl index 340f95e02..8c4a9398f 100644 --- a/test/PointerAccessChains/pointer_array_stride_16.cl +++ b/test/PointerAccessChains/pointer_array_stride_16.cl @@ -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; }; @@ -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 diff --git a/test/RewritePackedStructs/packed_struct.cl b/test/RewritePackedStructs/packed_struct.cl index f5b5e412a..3acb771f8 100644 --- a/test/RewritePackedStructs/packed_struct.cl +++ b/test/RewritePackedStructs/packed_struct.cl @@ -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; @@ -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