From 8fa0964caa5db11e94b70b56aa3122163b416016 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Wed, 7 Feb 2024 15:41:14 -0500 Subject: [PATCH] Rewrite UBO tests refs #1292 refs #1303 * Fixed a bug in remove unused arguments * attributes on parameters weren't handled parameters * rewrite tests to be more targeted at the actual feature --- lib/RemoveUnusedArguments.cpp | 9 +- test/UBO/array_stride_32.cl | 43 -------- test/UBO/array_stride_32.ll | 62 +++++++++++ test/UBO/can_dra_but_disabled.cl | 22 ---- test/UBO/can_dra_but_disabled.ll | 38 +++++++ test/UBO/can_dra_but_disabled_two_kernels.cl | 34 ------ test/UBO/can_dra_but_disabled_two_kernels.ll | 55 ++++++++++ test/UBO/cannot_dra.cl | 37 ------- test/UBO/cannot_dra.ll | 61 +++++++++++ test/UBO/char_ubo_struct.cl | 52 ---------- test/UBO/char_ubo_struct.ll | 64 ++++++++++++ test/UBO/char_ubo_struct_novec3.cl | 51 --------- test/UBO/char_ubo_struct_novec3.ll | 60 +++++++++++ test/UBO/clustered_pod_type_mutate.cl | 39 ------- test/UBO/clustered_pod_type_mutate.ll | 65 ++++++++++++ test/UBO/clustered_pod_ubo.cl | 46 --------- test/UBO/clustered_pod_ubo.ll | 67 ++++++++++++ test/UBO/constant_and_image.cl | 69 ------------- test/UBO/constant_and_image.ll | 103 +++++++++++++++++++ test/UBO/constant_wrapping.cl | 56 ---------- test/UBO/constant_wrapping.ll | 74 +++++++++++++ test/UBO/copy.cl | 28 ----- test/UBO/copy.ll | 47 +++++++++ test/UBO/copy_nested.cl | 40 ------- test/UBO/copy_nested.ll | 53 ++++++++++ test/UBO/extra_arg.cl | 36 ------- test/UBO/extra_arg.ll | 42 ++++++++ test/UBO/extra_args.cl | 45 -------- test/UBO/extra_args.ll | 50 +++++++++ test/UBO/global_wrapping.cl | 58 ----------- test/UBO/global_wrapping.ll | 75 ++++++++++++++ test/UBO/large_padding.cl | 41 -------- test/UBO/large_padding.ll | 54 ++++++++++ test/UBO/large_padding_std430.cl | 45 -------- test/UBO/large_padding_std430.ll | 59 +++++++++++ test/UBO/long_specialization_chain.cl | 74 ------------- test/UBO/long_specialization_chain.ll | 76 ++++++++++++++ test/UBO/max_ubo_size.cl | 45 -------- test/UBO/max_ubo_size.ll | 34 ++++++ test/UBO/mixed_inlining.cl | 35 ------- test/UBO/mixed_inlining.ll | 75 ++++++++++++++ test/UBO/multiple_ubo_args.cl | 47 --------- test/UBO/multiple_ubo_args.ll | 71 +++++++++++++ test/UBO/needs_inlined.cl | 18 ---- test/UBO/needs_inlined.ll | 55 ++++++++++ test/UBO/nested_padding.cl | 56 ---------- test/UBO/nested_padding.ll | 74 +++++++++++++ test/UBO/odd_size_padding.cl | 33 ------ test/UBO/odd_size_padding.ll | 45 ++++++++ test/UBO/test_cluster_pod_args.ll | 51 +++++++++ test/UBO/transform_local.cl | 71 ------------- test/UBO/transform_local.ll | 90 ++++++++++++++++ test/UBO/transform_padding.cl | 49 --------- test/UBO/transform_padding.ll | 62 +++++++++++ test/UBO/vec2_no_pad.cl | 42 -------- test/UBO/vec2_no_pad.ll | 55 ++++++++++ 56 files changed, 1724 insertions(+), 1214 deletions(-) delete mode 100644 test/UBO/array_stride_32.cl create mode 100644 test/UBO/array_stride_32.ll delete mode 100644 test/UBO/can_dra_but_disabled.cl create mode 100644 test/UBO/can_dra_but_disabled.ll delete mode 100644 test/UBO/can_dra_but_disabled_two_kernels.cl create mode 100644 test/UBO/can_dra_but_disabled_two_kernels.ll delete mode 100644 test/UBO/cannot_dra.cl create mode 100644 test/UBO/cannot_dra.ll delete mode 100644 test/UBO/char_ubo_struct.cl create mode 100644 test/UBO/char_ubo_struct.ll delete mode 100644 test/UBO/char_ubo_struct_novec3.cl create mode 100644 test/UBO/char_ubo_struct_novec3.ll delete mode 100644 test/UBO/clustered_pod_type_mutate.cl create mode 100644 test/UBO/clustered_pod_type_mutate.ll delete mode 100644 test/UBO/clustered_pod_ubo.cl create mode 100644 test/UBO/clustered_pod_ubo.ll delete mode 100644 test/UBO/constant_and_image.cl create mode 100644 test/UBO/constant_and_image.ll delete mode 100644 test/UBO/constant_wrapping.cl create mode 100644 test/UBO/constant_wrapping.ll delete mode 100644 test/UBO/copy.cl create mode 100644 test/UBO/copy.ll delete mode 100644 test/UBO/copy_nested.cl create mode 100644 test/UBO/copy_nested.ll delete mode 100644 test/UBO/extra_arg.cl create mode 100644 test/UBO/extra_arg.ll delete mode 100644 test/UBO/extra_args.cl create mode 100644 test/UBO/extra_args.ll delete mode 100644 test/UBO/global_wrapping.cl create mode 100644 test/UBO/global_wrapping.ll delete mode 100644 test/UBO/large_padding.cl create mode 100644 test/UBO/large_padding.ll delete mode 100644 test/UBO/large_padding_std430.cl create mode 100644 test/UBO/large_padding_std430.ll delete mode 100644 test/UBO/long_specialization_chain.cl create mode 100644 test/UBO/long_specialization_chain.ll delete mode 100644 test/UBO/max_ubo_size.cl create mode 100644 test/UBO/max_ubo_size.ll delete mode 100644 test/UBO/mixed_inlining.cl create mode 100644 test/UBO/mixed_inlining.ll delete mode 100644 test/UBO/multiple_ubo_args.cl create mode 100644 test/UBO/multiple_ubo_args.ll delete mode 100644 test/UBO/needs_inlined.cl create mode 100644 test/UBO/needs_inlined.ll delete mode 100644 test/UBO/nested_padding.cl create mode 100644 test/UBO/nested_padding.ll delete mode 100644 test/UBO/odd_size_padding.cl create mode 100644 test/UBO/odd_size_padding.ll create mode 100644 test/UBO/test_cluster_pod_args.ll delete mode 100644 test/UBO/transform_local.cl create mode 100644 test/UBO/transform_local.ll delete mode 100644 test/UBO/transform_padding.cl create mode 100644 test/UBO/transform_padding.ll delete mode 100644 test/UBO/vec2_no_pad.cl create mode 100644 test/UBO/vec2_no_pad.ll diff --git a/lib/RemoveUnusedArguments.cpp b/lib/RemoveUnusedArguments.cpp index 10863c0ba..82f22f367 100644 --- a/lib/RemoveUnusedArguments.cpp +++ b/lib/RemoveUnusedArguments.cpp @@ -74,11 +74,17 @@ void clspv::RemoveUnusedArguments::removeUnusedParameters( f->removeFromParent(); // Rebuild the type. + auto fn_attrs = f->getAttributes(); SmallVector arg_types; + uint32_t idx = 0; for (auto *arg : candidate.args) { if (arg) { arg_types.push_back(arg->getType()); + } else { + // Remove any parameter attributes for deleted args. + fn_attrs = fn_attrs.removeParamAttributes(M.getContext(), idx); } + idx++; } FunctionType *new_type = FunctionType::get(f->getReturnType(), arg_types, false); @@ -86,8 +92,7 @@ void clspv::RemoveUnusedArguments::removeUnusedParameters( // Insert the new function. Copy the calling convention, attributes and // metadata. auto inserted = - M.getOrInsertFunction(f->getName(), new_type, f->getAttributes()) - .getCallee(); + M.getOrInsertFunction(f->getName(), new_type, fn_attrs).getCallee(); Function *new_function = cast(inserted); new_function->setCallingConv(f->getCallingConv()); new_function->copyMetadata(f, 0); diff --git a/test/UBO/array_stride_32.cl b/test/UBO/array_stride_32.cl deleted file mode 100644 index 269b17785..000000000 --- a/test/UBO/array_stride_32.cl +++ /dev/null @@ -1,43 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// Checking for correct offsets and array strides. - -typedef struct { - int x; - int y __attribute((aligned(16))); -} s; - -__kernel void foo(__global s* data, __constant s* c) { - unsigned gid = get_global_id(0); - data[gid].x = c[gid].x; - data[gid].y = c[gid].y; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 4 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 16 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 20 -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 32 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 -// CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK: [[s]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] -// CHECK: [[runtime]] = OpTypeRuntimeArray [[s]] -// CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] -// CHECK: [[int_2048:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2048 -// CHECK: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[s]] [[int_2048]] -// CHECK: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c]] = OpVariable [[c_ptr]] Uniform diff --git a/test/UBO/array_stride_32.ll b/test/UBO/array_stride_32.ll new file mode 100644 index 000000000..d7c24d57c --- /dev/null +++ b/test/UBO/array_stride_32.ll @@ -0,0 +1,62 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file=%t.spv -int8=0 --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 4 +; CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 16 +; CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 20 +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 32 +; CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 +; CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK: [[s]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] +; CHECK: [[runtime]] = OpTypeRuntimeArray [[s]] +; CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] +; CHECK: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] +; CHECK: [[int_2048:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2048 +; CHECK: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[s]] [[int_2048]] +; CHECK: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] +; CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[c]] = OpVariable [[c_ptr]] Uniform + +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" + +%struct.s = type { i32, [12 x i8], i32, [12 x i8] } + +@__spirv_GlobalInvocationId = local_unnamed_addr addrspace(5) global <3 x i32> zeroinitializer +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !13 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.s] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [2048 x %struct.s] } zeroinitializer) + %2 = getelementptr <3 x i32>, ptr addrspace(5) @__spirv_GlobalInvocationId, i32 0, i32 0 + %3 = load i32, ptr addrspace(5) %2, align 16 + %4 = getelementptr { [2048 x %struct.s] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 0 + %5 = load i32, ptr addrspace(2) %4, align 16 + %6 = getelementptr { [0 x %struct.s] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 0 + store i32 %5, ptr addrspace(1) %6, align 16 + %7 = getelementptr { [2048 x %struct.s] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 2 + %8 = load i32, ptr addrspace(2) %7, align 16 + %9 = getelementptr { [0 x %struct.s] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 2 + store i32 %8, ptr addrspace(1) %9, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.s] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [2048 x %struct.s] }) + +!13 = !{i32 2} + diff --git a/test/UBO/can_dra_but_disabled.cl b/test/UBO/can_dra_but_disabled.cl deleted file mode 100644 index a6cb65fa5..000000000 --- a/test/UBO/can_dra_but_disabled.cl +++ /dev/null @@ -1,22 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -no-dra -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -__attribute__((noinline)) -int4 bar(constant int4* data) { return data[0]; } - -kernel void foo(global int4* out, constant int4* in) { - *out = bar(in); -} - -// CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] -// CHECK-DAG: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK-DAG: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK: [[foo]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar:%[a-zA-Z0-9_]+]] -// CHECK: [[bar]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int0]] -// CHECK: OpLoad {{.*}} [[gep]] diff --git a/test/UBO/can_dra_but_disabled.ll b/test/UBO/can_dra_but_disabled.ll new file mode 100644 index 000000000..7e48c056f --- /dev/null +++ b/test/UBO/can_dra_but_disabled.ll @@ -0,0 +1,38 @@ +; RUN: clspv-opt %s -o %t.ll -constant-args-ubo --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} void @foo +; CHECK: call <4 x i32> [[bar:@bar[a-zA-Z0-9_]+]] +; CHECK: define {{.*}} <4 x i32> [[bar]]() +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 0 +; CHECK: load <4 x i32>, ptr addrspace(2) [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %data) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %data, align 16 + ret <4 x i32> %0 +} + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !16 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %3) + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!16 = !{i32 2} + diff --git a/test/UBO/can_dra_but_disabled_two_kernels.cl b/test/UBO/can_dra_but_disabled_two_kernels.cl deleted file mode 100644 index f82d3fef8..000000000 --- a/test/UBO/can_dra_but_disabled_two_kernels.cl +++ /dev/null @@ -1,34 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -no-dra -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -__attribute__((noinline)) -int4 bar(constant int4* data) { return data[0]; } - -kernel void k1(global int4* out, constant int4* in) { - *out = bar(in); -} - -kernel void k2(global int4* out, constant int4* in) { - *out = bar(in); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: OpEntryPoint GLCompute [[k2:%[a-zA-Z0-9_]+]] -// CHECK-DAG: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK-DAG: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK: [[k1]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar1:%[a-zA-Z0-9_]+]] -// CHECK: [[k2]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar2:%[a-zA-Z0-9_]+]] -// CHECK: [[bar2]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int0]] -// CHECK: OpLoad {{.*}} [[gep]] -// CHECK: [[bar1]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int0]] -// CHECK: OpLoad {{.*}} [[gep]] - diff --git a/test/UBO/can_dra_but_disabled_two_kernels.ll b/test/UBO/can_dra_but_disabled_two_kernels.ll new file mode 100644 index 000000000..7a331d36e --- /dev/null +++ b/test/UBO/can_dra_but_disabled_two_kernels.ll @@ -0,0 +1,55 @@ +; RUN: clspv-opt %s -o %t.ll -constant-args-ubo --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} @k1 +; CHECK: call {{.*}} [[bar1:@bar[a-zA-Z0-9_.]+]]() +; CHECK: define {{.*}} @k2 +; CHECK: call {{.*}} [[bar2:@bar[a-zA-Z0-9_.]+]]() +; CHECK-DAG: define {{.*}} [[bar1]]() +; CHECK-DAG: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK-DAG: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 0 +; CHECK-DAG: load <4 x i32>, ptr addrspace(2) [[gep]] +; CHECK-DAG: define {{.*}} [[bar2]]() +; CHECK-DAG: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK-DAG: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 0 +; CHECK-DAG: load <4 x i32>, ptr addrspace(2) [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %data) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %data, align 16 + ret <4 x i32> %0 +} + +define dso_local spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !17 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %3) + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +define dso_local spir_kernel void @k2(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !17 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %3) + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!17 = !{i32 2} + diff --git a/test/UBO/cannot_dra.cl b/test/UBO/cannot_dra.cl deleted file mode 100644 index 1c41445f7..000000000 --- a/test/UBO/cannot_dra.cl +++ /dev/null @@ -1,37 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -// TODO(#1292) -// XFAIL: * - -__attribute__((noinline)) -int4 bar(constant int4* data) { return data[0]; } - -kernel void k1(global int4* out, constant int4* in) { - *out = bar(in); -} - -kernel void k2(global int4* out, constant int4* in) { - *out = bar(in + 1); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: OpEntryPoint GLCompute [[k2:%[a-zA-Z0-9_]+]] -// CHECK-DAG: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK-DAG: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[int1:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 1 -// CHECK: [[k1]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar1:%[a-zA-Z0-9_]+]] -// CHECK: [[k2]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar2:%[a-zA-Z0-9_]+]] -// CHECK: [[bar2]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int1]] -// CHECK: OpLoad {{.*}} [[gep]] -// CHECK: [[bar1]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int0]] -// CHECK: OpLoad {{.*}} [[gep]] diff --git a/test/UBO/cannot_dra.ll b/test/UBO/cannot_dra.ll new file mode 100644 index 000000000..ac50573c3 --- /dev/null +++ b/test/UBO/cannot_dra.ll @@ -0,0 +1,61 @@ +; RUN: clspv-opt %s -o %t.ll -constant-args-ubo --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} @k1 +; CHECK: call {{.*}} [[bar1:@bar[a-zA-Z0-9_.]+]] +; CHECK: define {{.*}} @k2 +; CHECK: call {{.*}} [[bar2:@bar[a-zA-Z0-9_.]+]] +; CHECK-DAG: define {{.*}} [[bar1]]() +; CHECK-DAG: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK-DAG: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 0 +; CHECK-DAG: load <4 x i32>, ptr addrspace(2) [[gep]] +; CHECK-DAG: define {{.*}} [[bar2]]() +; CHECK-DAG: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK-DAG: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 1 +; CHECK-DAG: load <4 x i32>, ptr addrspace(2) [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %data) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %data, align 16 + ret <4 x i32> %0 +} + +define dso_local spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !17 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %3) + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +define spir_kernel void @k2(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !17 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = call ptr addrspace(9) @_Z14clspv.resource.2(i32 -1, i32 2, i32 5, i32 2, i32 2, i32 0, { { i32 } } zeroinitializer) + %4 = getelementptr { { i32 } }, ptr addrspace(9) %3, i32 0, i32 0 + %5 = load { i32 }, ptr addrspace(9) %4, align 4 + %x = extractvalue { i32 } %5, 0 + %6 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 1 + %call.i = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %6) + store <4 x i32> %call.i, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +declare ptr addrspace(9) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { { i32 } }) + +!17 = !{i32 2} + diff --git a/test/UBO/char_ubo_struct.cl b/test/UBO/char_ubo_struct.cl deleted file mode 100644 index bb7233a16..000000000 --- a/test/UBO/char_ubo_struct.cl +++ /dev/null @@ -1,52 +0,0 @@ -// RUN: clspv %target -int8 -constant-args-ubo -inline-entry-points %s -o %t.spv -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: clspv-reflection %t.spv -o %t.map -// RUN: FileCheck -check-prefix=MAP %s < %t.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// TODO(#1292) -// XFAIL: * - -typedef struct { - char a; - char2 b; - char3 c; - char4 d; - int pad; // necessary to get up to 16 byte size -} S; - -kernel void foo(global S* out, constant S* in) { - out->d = in->d; -} - -// MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 2 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 4 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 8 -// CHECK-DAG: OpMemberDecorate [[s]] 4 Offset 12 -// CHECK: OpDecorate [[rta:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: OpDecorate [[out:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[out]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[in]] NonWritable -// CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 -// CHECK-DAG: [[char2:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 2 -// CHECK-DAG: [[char3:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 3 -// CHECK-DAG: [[char4:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 4 -// CHECK: [[s]] = OpTypeStruct [[char]] [[char2]] [[char3]] [[char4]] [[int]] -// CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_4096]] -// CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] -// CHECK-DAG: [[rta]] = OpTypeRuntimeArray [[s]] -// CHECK-DAG: [[ssbo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[rta]] -// CHECK-DAG: [[ssbo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[ssbo_block]] -// CHECK-DAG: [[out]] = OpVariable [[ssbo_ptr]] StorageBuffer -// CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform diff --git a/test/UBO/char_ubo_struct.ll b/test/UBO/char_ubo_struct.ll new file mode 100644 index 000000000..4aa19513c --- /dev/null +++ b/test/UBO/char_ubo_struct.ll @@ -0,0 +1,64 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 2 +; CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 4 +; CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 8 +; CHECK-DAG: OpMemberDecorate [[s]] 4 Offset 12 +; CHECK: OpDecorate [[rta:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: OpDecorate [[out:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[out]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[in]] NonWritable +; CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 +; CHECK-DAG: [[char2:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 2 +; CHECK-DAG: [[char3:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 3 +; CHECK-DAG: [[char4:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 4 +; CHECK: [[s]] = OpTypeStruct [[char]] [[char2]] [[char3]] [[char4]] [[int]] +; CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_4096]] +; CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] +; CHECK-DAG: [[rta]] = OpTypeRuntimeArray [[s]] +; CHECK-DAG: [[ssbo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[rta]] +; CHECK-DAG: [[ssbo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[ssbo_block]] +; CHECK-DAG: [[out]] = OpVariable [[ssbo_ptr]] StorageBuffer +; CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform + +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" + +%struct.S = type { i8, <2 x i8>, <3 x i8>, <4 x i8>, i32 } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 4 %out, ptr addrspace(2) nocapture readonly align 4 %in) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.S] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x %struct.S] } zeroinitializer) + %5 = getelementptr { [4096 x %struct.S] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0, i32 3 + %6 = load <4 x i8>, ptr addrspace(2) %5, align 4 + %7 = getelementptr { [0 x %struct.S] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0, i32 3 + store <4 x i8> %6, ptr addrspace(1) %7, align 4 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.S] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x %struct.S] }) + +declare ptr addrspace(9) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { { i32 } }) + +!14 = !{i32 2} + diff --git a/test/UBO/char_ubo_struct_novec3.cl b/test/UBO/char_ubo_struct_novec3.cl deleted file mode 100644 index c3382b194..000000000 --- a/test/UBO/char_ubo_struct_novec3.cl +++ /dev/null @@ -1,51 +0,0 @@ -// RUN: clspv %target -int8 -constant-args-ubo -inline-entry-points %s -o %t.spv -vec3-to-vec4 -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: clspv-reflection %t.spv -o %t.map -// RUN: FileCheck -check-prefix=MAP %s < %t.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// TODO(#1292) -// XFAIL: * - -typedef struct { - char a; - char2 b; - char3 c; - char4 d; - int pad; // necessary to get up to 16 byte size -} S; - -kernel void foo(global S* out, constant S* in) { - out->d = in->d; -} - -// MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 2 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 4 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 8 -// CHECK-DAG: OpMemberDecorate [[s]] 4 Offset 12 -// CHECK: OpDecorate [[rta:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: OpDecorate [[out:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[out]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[in]] NonWritable -// CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 -// CHECK-DAG: [[char2:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 2 -// CHECK-DAG: [[char4:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 4 -// CHECK: [[s]] = OpTypeStruct [[char]] [[char2]] [[char4]] [[char4]] [[int]] -// CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_4096]] -// CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] -// CHECK-DAG: [[rta]] = OpTypeRuntimeArray [[s]] -// CHECK-DAG: [[ssbo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[rta]] -// CHECK-DAG: [[ssbo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[ssbo_block]] -// CHECK-DAG: [[out]] = OpVariable [[ssbo_ptr]] StorageBuffer -// CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform diff --git a/test/UBO/char_ubo_struct_novec3.ll b/test/UBO/char_ubo_struct_novec3.ll new file mode 100644 index 000000000..690a4bf57 --- /dev/null +++ b/test/UBO/char_ubo_struct_novec3.ll @@ -0,0 +1,60 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 2 +; CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 4 +; CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 8 +; CHECK-DAG: OpMemberDecorate [[s]] 4 Offset 12 +; CHECK: OpDecorate [[rta:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: OpDecorate [[out:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[out]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[in]] NonWritable +; CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 +; CHECK-DAG: [[char2:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 2 +; CHECK-DAG: [[char4:%[0-9a-zA-Z_]+]] = OpTypeVector [[char]] 4 +; CHECK: [[s]] = OpTypeStruct [[char]] [[char2]] [[char4]] [[char4]] [[int]] +; CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_4096]] +; CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] +; CHECK-DAG: [[rta]] = OpTypeRuntimeArray [[s]] +; CHECK-DAG: [[ssbo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[rta]] +; CHECK-DAG: [[ssbo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[ssbo_block]] +; CHECK-DAG: [[out]] = OpVariable [[ssbo_ptr]] StorageBuffer +; CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 4 %out, ptr addrspace(2) nocapture readonly align 4 %in) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x { i8, <2 x i8>, <4 x i8>, <4 x i8>, i32 }] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x { i8, <2 x i8>, <4 x i8>, <4 x i8>, i32 }] } zeroinitializer) + %5 = getelementptr { [4096 x { i8, <2 x i8>, <4 x i8>, <4 x i8>, i32 }] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0, i32 3 + %6 = load <4 x i8>, ptr addrspace(2) %5, align 4 + %7 = getelementptr { [0 x { i8, <2 x i8>, <4 x i8>, <4 x i8>, i32 }] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0, i32 3 + store <4 x i8> %6, ptr addrspace(1) %7, align 4 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x { i8, <2 x i8>, <4 x i8>, <4 x i8>, i32 }] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x { i8, <2 x i8>, <4 x i8>, <4 x i8>, i32 }] }) + + +!14 = !{i32 2} + diff --git a/test/UBO/clustered_pod_type_mutate.cl b/test/UBO/clustered_pod_type_mutate.cl deleted file mode 100644 index 2e5e7951e..000000000 --- a/test/UBO/clustered_pod_type_mutate.cl +++ /dev/null @@ -1,39 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -pod-ubo -cluster-pod-kernel-args -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: clspv-reflection %t.spv -o %t.map -// RUN: FileCheck --check-prefix=MAP %s < %t.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -typedef struct { - int a __attribute__((aligned(16))); -} S; - -kernel void foo(global int4* out, int a, S s) { - out->x = a + s.a; -} - -// MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,a,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,pod_ubo,argSize,4 -// MAP-NEXT: kernel,foo,arg,s,argOrdinal,2,descriptorSet,0,binding,1,offset,16,argKind,pod_ubo,argSize,16 - -// CHECK: OpMemberDecorate -// CHECK: OpMemberDecorate [[S:%[a-zA-Z0-9_]+]] 0 Offset 0 -// CHECK: OpMemberDecorate [[S]] 1 Offset 4 -// CHECK: OpMemberDecorate [[cluster:%[a-zA-Z0-9_]+]] 0 Offset 0 -// CHECK: OpMemberDecorate [[cluster]] 1 Offset 4 -// CHECK: OpMemberDecorate [[cluster]] 2 Offset 8 -// CHECK: OpMemberDecorate [[cluster]] 3 Offset 12 -// CHECK: OpMemberDecorate [[cluster]] 4 Offset 16 -// CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK: [[char:%[a-zA-Z0-9_]+]] = OpTypeInt 8 0 -// CHECK: [[S]] = OpTypeStruct [[int]] [[char]] -// CHECK: [[cluster]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] [[S]] -// CHECK: [[pod_var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[pod_var]] -// CHECK: [[ld:%[a-zA-Z0-9_]+]] = OpLoad [[cluster]] [[gep]] -// CHECK: [[a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[ld]] 0 -// CHECK: [[s:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[S]] [[ld]] 4 -// CHECK: [[s_a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[s]] 0 -// CHECK: OpIAdd [[int]] [[s_a]] [[a]] - diff --git a/test/UBO/clustered_pod_type_mutate.ll b/test/UBO/clustered_pod_type_mutate.ll new file mode 100644 index 000000000..ec19acc7e --- /dev/null +++ b/test/UBO/clustered_pod_type_mutate.ll @@ -0,0 +1,65 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,a,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,pod_ubo,argSize,4 +; MAP-NEXT: kernel,foo,arg,s,argOrdinal,2,descriptorSet,0,binding,1,offset,16,argKind,pod_ubo,argSize,16 + +; CHECK: OpMemberDecorate +; CHECK: OpMemberDecorate [[S:%[a-zA-Z0-9_]+]] 0 Offset 0 +; CHECK: OpMemberDecorate [[S]] 1 Offset 4 +; CHECK: OpMemberDecorate [[cluster:%[a-zA-Z0-9_]+]] 0 Offset 0 +; CHECK: OpMemberDecorate [[cluster]] 1 Offset 4 +; CHECK: OpMemberDecorate [[cluster]] 2 Offset 8 +; CHECK: OpMemberDecorate [[cluster]] 3 Offset 12 +; CHECK: OpMemberDecorate [[cluster]] 4 Offset 16 +; CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 +; CHECK: [[char:%[a-zA-Z0-9_]+]] = OpTypeInt 8 0 +; CHECK: [[S]] = OpTypeStruct [[int]] [[char]] +; CHECK: [[cluster]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] [[S]] +; CHECK: [[pod_var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform +; CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[pod_var]] +; CHECK: [[ld:%[a-zA-Z0-9_]+]] = OpLoad [[cluster]] [[gep]] +; CHECK: [[a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[ld]] 0 +; CHECK: [[s:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[S]] [[ld]] 4 +; CHECK: [[s_a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[s]] 0 +; CHECK: OpIAdd [[int]] [[s_a]] [[a]] + +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" + +%struct.S = type { i32, [12 x i8] } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture align 16 %out, { i32, i32, i32, i32, %struct.S } %podargs) !clspv.pod_args_impl !8 !kernel_arg_map !15 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(6) @_Z14clspv.resource.1(i32 0, i32 1, i32 4, i32 1, i32 1, i32 0, { { i32, i32, i32, i32, %struct.S } } zeroinitializer) + %3 = getelementptr { { i32, i32, i32, i32, %struct.S } }, ptr addrspace(6) %2, i32 0, i32 0 + %4 = load { i32, i32, i32, i32, %struct.S }, ptr addrspace(6) %3, align 4 + %a = extractvalue { i32, i32, i32, i32, %struct.S } %4, 0 + %s = extractvalue { i32, i32, i32, i32, %struct.S } %4, 4 + %s.elt = extractvalue %struct.S %s, 0 + %add.i = add nsw i32 %s.elt, %a + %5 = load <4 x i32>, ptr addrspace(1) %1, align 16 + %6 = insertelement <4 x i32> %5, i32 %add.i, i64 0 + store <4 x i32> %6, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(6) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { { i32, i32, i32, i32, %struct.S } }) + +!8 = !{i32 1} +!15 = !{!16, !17, !18} +!16 = !{!"out", i32 0, i32 0, i32 0, i32 0, !"buffer"} +!17 = !{!"a", i32 1, i32 1, i32 0, i32 4, !"pod_ubo"} +!18 = !{!"s", i32 2, i32 1, i32 16, i32 16, !"pod_ubo"} + diff --git a/test/UBO/clustered_pod_ubo.cl b/test/UBO/clustered_pod_ubo.cl deleted file mode 100644 index 669af2b42..000000000 --- a/test/UBO/clustered_pod_ubo.cl +++ /dev/null @@ -1,46 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -pod-ubo -cluster-pod-kernel-args -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: clspv-reflection %t.spv -o %t.map -// RUN: FileCheck --check-prefix=MAP %s < %t.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv -// RUN: clspv %target %s -o %t2.spv -pod-ubo -cluster-pod-kernel-args -int8=0 -// RUN: spirv-dis %t2.spv -o %t2.spvasm -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t2.spv -o %t2.map -// RUN: FileCheck --check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t2.spv - -typedef struct { - int a, b, c, d; -} S; - -kernel void foo(global int4* out, int a, S s) { - out->x = a + s.a; -} - -// MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,a,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,pod_ubo,argSize,4 -// MAP-NEXT: kernel,foo,arg,s,argOrdinal,2,descriptorSet,0,binding,1,offset,16,argKind,pod_ubo,argSize,16 - -// CHECK: OpMemberDecorate -// CHECK: OpMemberDecorate [[S:%[a-zA-Z0-9_]+]] 0 Offset 0 -// CHECK: OpMemberDecorate [[S]] 1 Offset 4 -// CHECK: OpMemberDecorate [[S]] 2 Offset 8 -// CHECK: OpMemberDecorate [[S]] 3 Offset 12 -// CHECK: OpMemberDecorate [[cluster:%[a-zA-Z0-9_]+]] 0 Offset 0 -// CHECK: OpMemberDecorate [[cluster]] 1 Offset 4 -// CHECK: OpMemberDecorate [[cluster]] 2 Offset 8 -// CHECK: OpMemberDecorate [[cluster]] 3 Offset 12 -// CHECK: OpMemberDecorate [[cluster]] 4 Offset 16 -// CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK-NOT: OpTypeInt 8 0 -// CHECK: [[S]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] -// CHECK: [[cluster]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] [[S]] -// CHECK: [[pod_var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[pod_var]] -// CHECK: [[ld:%[a-zA-Z0-9_]+]] = OpLoad [[cluster]] [[gep]] -// CHECK: [[a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[ld]] 0 -// CHECK: [[s:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[S]] [[ld]] 4 -// CHECK: [[s_a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[s]] 0 -// CHECK: OpIAdd [[int]] [[s_a]] [[a]] diff --git a/test/UBO/clustered_pod_ubo.ll b/test/UBO/clustered_pod_ubo.ll new file mode 100644 index 000000000..68320f5b0 --- /dev/null +++ b/test/UBO/clustered_pod_ubo.ll @@ -0,0 +1,67 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,a,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,pod_ubo,argSize,4 +; MAP-NEXT: kernel,foo,arg,s,argOrdinal,2,descriptorSet,0,binding,1,offset,16,argKind,pod_ubo,argSize,16 + +; CHECK: OpMemberDecorate +; CHECK: OpMemberDecorate [[S:%[a-zA-Z0-9_]+]] 0 Offset 0 +; CHECK: OpMemberDecorate [[S]] 1 Offset 4 +; CHECK: OpMemberDecorate [[S]] 2 Offset 8 +; CHECK: OpMemberDecorate [[S]] 3 Offset 12 +; CHECK: OpMemberDecorate [[cluster:%[a-zA-Z0-9_]+]] 0 Offset 0 +; CHECK: OpMemberDecorate [[cluster]] 1 Offset 4 +; CHECK: OpMemberDecorate [[cluster]] 2 Offset 8 +; CHECK: OpMemberDecorate [[cluster]] 3 Offset 12 +; CHECK: OpMemberDecorate [[cluster]] 4 Offset 16 +; CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 +; CHECK-NOT: OpTypeInt 8 0 +; CHECK: [[S]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] +; CHECK: [[cluster]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] [[S]] +; CHECK: [[pod_var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform +; CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[pod_var]] +; CHECK: [[ld:%[a-zA-Z0-9_]+]] = OpLoad [[cluster]] [[gep]] +; CHECK: [[a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[ld]] 0 +; CHECK: [[s:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[S]] [[ld]] 4 +; CHECK: [[s_a:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] [[s]] 0 +; CHECK: OpIAdd [[int]] [[s_a]] [[a]] + +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" + +%struct.S = type { i32, i32, i32, i32 } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture align 16 %out, { i32, i32, i32, i32, %struct.S } %podargs) !clspv.pod_args_impl !8 !kernel_arg_map !15 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(6) @_Z14clspv.resource.1(i32 0, i32 1, i32 4, i32 1, i32 1, i32 0, { { i32, i32, i32, i32, %struct.S } } zeroinitializer) + %3 = getelementptr { { i32, i32, i32, i32, %struct.S } }, ptr addrspace(6) %2, i32 0, i32 0 + %4 = load { i32, i32, i32, i32, %struct.S }, ptr addrspace(6) %3, align 4 + %a = extractvalue { i32, i32, i32, i32, %struct.S } %4, 0 + %s = extractvalue { i32, i32, i32, i32, %struct.S } %4, 4 + %s.elt = extractvalue %struct.S %s, 0 + %add.i = add nsw i32 %s.elt, %a + %5 = load <4 x i32>, ptr addrspace(1) %1, align 16 + %6 = insertelement <4 x i32> %5, i32 %add.i, i64 0 + store <4 x i32> %6, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(6) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { { i32, i32, i32, i32, %struct.S } }) + +!8 = !{i32 1} +!15 = !{!16, !17, !18} +!16 = !{!"out", i32 0, i32 0, i32 0, i32 0, !"buffer"} +!17 = !{!"a", i32 1, i32 1, i32 0, i32 4, !"pod_ubo"} +!18 = !{!"s", i32 2, i32 1, i32 16, i32 16, !"pod_ubo"} + diff --git a/test/UBO/constant_and_image.cl b/test/UBO/constant_and_image.cl deleted file mode 100644 index 6a5d01f50..000000000 --- a/test/UBO/constant_and_image.cl +++ /dev/null @@ -1,69 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -cluster-pod-kernel-args=0 -pod-ubo -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -kernel void foo(read_only image2d_t i, sampler_t s, constant float4* offset, float2 c, global float4* data) { - *data = read_imagef(i, s, c) + *offset; -} - -// MAP: kernel,foo,arg,i,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,ro_image -// MAP-NEXT: kernel,foo,arg,s,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,sampler -// MAP-NEXT: kernel,foo,arg,offset,argOrdinal,2,descriptorSet,0,binding,2,offset,0,argKind,buffer_ubo -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,3,descriptorSet,0,binding,3,offset,0,argKind,pod_ubo,argSize,8 -// MAP-NEXT: kernel,foo,arg,data,argOrdinal,4,descriptorSet,0,binding,4,offset,0,argKind,buffer - -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: OpDecorate [[image_var:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[image_var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[sampler_var:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[sampler_var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[offset_var:%[0-9a-zA-Z_]+]] Binding 2 -// CHECK-DAG: OpDecorate [[offset_var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[offset_var]] NonWritable -// CHECK-DAG: OpDecorate [[c_var:%[0-9a-zA-Z_]+]] Binding 3 -// CHECK-DAG: OpDecorate [[c_var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[data_var:%[0-9a-zA-Z_]+]] Binding 4 -// CHECK-DAG: OpDecorate [[data_var]] DescriptorSet 0 -// CHECK-NOT: OpExtension -// CHECK-DAG: [[float:%[0-9a-zA-Z_]+]] = OpTypeFloat 32 -// CHECK-DAG: [[image:%[0-9a-zA-Z_]+]] = OpTypeImage [[float]] 2D 0 0 0 1 Unknown -// CHECK-DAG: [[sampled_image:%[0-9a-zA-Z_]+]] = OpTypeSampledImage [[image]] -// CHECK-DAG: [[image_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer UniformConstant [[image]] -// CHECK-DAG: [[sampler:%[0-9a-zA-Z_]+]] = OpTypeSampler -// CHECK-DAG: [[sampler_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer UniformConstant [[sampler]] -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[v4float:%[0-9a-zA-Z_]+]] = OpTypeVector [[float]] 4 -// CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[v4float]] [[int_4096]] -// CHECK-DAG: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK-DAG: [[offset_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] -// CHECK-DAG: [[v2float:%[0-9a-zA-Z_]+]] = OpTypeVector [[float]] 2 -// CHECK-DAG: [[struct_v2float:%[0-9a-zA-Z_]+]] = OpTypeStruct [[v2float]] -// CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct_v2float]] -// CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[v4float]] -// CHECK-DAG: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] -// CHECK-DAG: [[ptr_uniform_v4float:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[v4float]] -// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[ptr_uniform_v2float:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[v2float]] -// CHECK-DAG: [[ptr_storagebuffer_v4float:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[v4float]] -// CHECK-DAG: [[float_zero:%[0-9a-zA-Z_]+]] = OpConstant [[float]] 0 -// CHECK: [[image_var]] = OpVariable [[image_ptr]] UniformConstant -// CHECK: [[sampler_var]] = OpVariable [[sampler_ptr]] UniformConstant -// CHECK: [[offset_var]] = OpVariable [[offset_ptr]] Uniform -// CHECK: [[c_var]] = OpVariable [[c_ptr]] Uniform -// CHECK: [[data_var]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[load_image:%[0-9a-zA-Z_]+]] = OpLoad [[image]] [[image_var]] -// CHECK: [[load_sampler:%[0-9a-zA-Z_]+]] = OpLoad [[sampler]] [[sampler_var]] -// CHECK: [[offset_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_uniform_v4float]] [[offset_var]] [[zero]] [[zero]] -// CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_uniform_v2float]] [[c_var]] [[zero]] -// CHECK: [[load_c:%[0-9a-zA-Z_]+]] = OpLoad [[v2float]] [[c_gep]] -// CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_storagebuffer_v4float]] [[data_var]] [[zero]] [[zero]] -// CHECK: [[sampled:%[0-9a-zA-Z_]+]] = OpSampledImage [[sampled_image]] [[load_image]] [[load_sampler]] -// CHECK: [[sample:%[0-9a-zA-Z_]+]] = OpImageSampleExplicitLod [[v4float]] [[sampled]] [[load_c]] Lod [[float_zero]] -// CHECK: [[offset_load:%[0-9a-zA-Z_]+]] = OpLoad [[v4float]] [[offset_gep]] -// CHECK: [[add:%[0-9a-zA-Z_]+]] = OpFAdd [[v4float]] [[sample]] [[offset_load]] -// CHECK: OpStore [[data_gep]] [[add]] diff --git a/test/UBO/constant_and_image.ll b/test/UBO/constant_and_image.ll new file mode 100644 index 000000000..50f25047a --- /dev/null +++ b/test/UBO/constant_and_image.ll @@ -0,0 +1,103 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,i,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,ro_image +; MAP-NEXT: kernel,foo,arg,s,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,sampler +; MAP-NEXT: kernel,foo,arg,offset,argOrdinal,2,descriptorSet,0,binding,2,offset,0,argKind,buffer_ubo +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,3,descriptorSet,0,binding,3,offset,0,argKind,pod_ubo,argSize,8 +; MAP-NEXT: kernel,foo,arg,data,argOrdinal,4,descriptorSet,0,binding,4,offset,0,argKind,buffer + +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: OpDecorate [[image_var:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[image_var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[sampler_var:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[sampler_var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[offset_var:%[0-9a-zA-Z_]+]] Binding 2 +; CHECK-DAG: OpDecorate [[offset_var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[offset_var]] NonWritable +; CHECK-DAG: OpDecorate [[c_var:%[0-9a-zA-Z_]+]] Binding 3 +; CHECK-DAG: OpDecorate [[c_var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[data_var:%[0-9a-zA-Z_]+]] Binding 4 +; CHECK-DAG: OpDecorate [[data_var]] DescriptorSet 0 +; CHECK-NOT: OpExtension +; CHECK-DAG: [[float:%[0-9a-zA-Z_]+]] = OpTypeFloat 32 +; CHECK-DAG: [[image:%[0-9a-zA-Z_]+]] = OpTypeImage [[float]] 2D 0 0 0 1 Unknown +; CHECK-DAG: [[sampled_image:%[0-9a-zA-Z_]+]] = OpTypeSampledImage [[image]] +; CHECK-DAG: [[image_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer UniformConstant [[image]] +; CHECK-DAG: [[sampler:%[0-9a-zA-Z_]+]] = OpTypeSampler +; CHECK-DAG: [[sampler_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer UniformConstant [[sampler]] +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[v4float:%[0-9a-zA-Z_]+]] = OpTypeVector [[float]] 4 +; CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[v4float]] [[int_4096]] +; CHECK-DAG: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK-DAG: [[offset_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] +; CHECK-DAG: [[v2float:%[0-9a-zA-Z_]+]] = OpTypeVector [[float]] 2 +; CHECK-DAG: [[struct_v2float:%[0-9a-zA-Z_]+]] = OpTypeStruct [[v2float]] +; CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct_v2float]] +; CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[v4float]] +; CHECK-DAG: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] +; CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] +; CHECK-DAG: [[ptr_uniform_v4float:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[v4float]] +; CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK-DAG: [[ptr_uniform_v2float:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[v2float]] +; CHECK-DAG: [[ptr_storagebuffer_v4float:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[v4float]] +; CHECK-DAG: [[float_zero:%[0-9a-zA-Z_]+]] = OpConstant [[float]] 0 +; CHECK: [[image_var]] = OpVariable [[image_ptr]] UniformConstant +; CHECK: [[sampler_var]] = OpVariable [[sampler_ptr]] UniformConstant +; CHECK: [[offset_var]] = OpVariable [[offset_ptr]] Uniform +; CHECK: [[c_var]] = OpVariable [[c_ptr]] Uniform +; CHECK: [[data_var]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[load_image:%[0-9a-zA-Z_]+]] = OpLoad [[image]] [[image_var]] +; CHECK: [[load_sampler:%[0-9a-zA-Z_]+]] = OpLoad [[sampler]] [[sampler_var]] +; CHECK: [[offset_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_uniform_v4float]] [[offset_var]] [[zero]] [[zero]] +; CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_uniform_v2float]] [[c_var]] [[zero]] +; CHECK: [[load_c:%[0-9a-zA-Z_]+]] = OpLoad [[v2float]] [[c_gep]] +; CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_storagebuffer_v4float]] [[data_var]] [[zero]] [[zero]] +; CHECK: [[sampled:%[0-9a-zA-Z_]+]] = OpSampledImage [[sampled_image]] [[load_image]] [[load_sampler]] +; CHECK: [[sample:%[0-9a-zA-Z_]+]] = OpImageSampleExplicitLod [[v4float]] [[sampled]] [[load_c]] Lod [[float_zero]] +; CHECK: [[offset_load:%[0-9a-zA-Z_]+]] = OpLoad [[v4float]] [[offset_gep]] +; CHECK: [[add:%[0-9a-zA-Z_]+]] = OpFAdd [[v4float]] [[sample]] [[offset_load]] +; CHECK: OpStore [[data_gep]] [[add]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +declare spir_func <4 x float> @_Z11read_imagef30ocl_image2d_ro_t.float.sampled11ocl_samplerDv2_f(target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0), target("spirv.Sampler"), <2 x float>) + +define spir_kernel void @foo(target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0) %i, target("spirv.Sampler") %s, ptr addrspace(2) nocapture readonly align 16 %offset, <2 x float> %c, ptr addrspace(1) nocapture writeonly align 16 %data) !clspv.pod_args_impl !8 { +entry: + %0 = call target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0) @_Z14clspv.resource.0(i32 0, i32 0, i32 6, i32 0, i32 0, i32 0, target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0) undef) + %1 = call target("spirv.Sampler") @_Z14clspv.resource.1(i32 0, i32 1, i32 8, i32 1, i32 1, i32 0, target("spirv.Sampler") zeroinitializer) + %2 = call ptr addrspace(2) @_Z14clspv.resource.2(i32 0, i32 2, i32 1, i32 2, i32 2, i32 0, { [4096 x <4 x float>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x float>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = call ptr addrspace(6) @_Z14clspv.resource.3(i32 0, i32 3, i32 4, i32 3, i32 3, i32 0, { <2 x float> } zeroinitializer) + %5 = getelementptr { <2 x float> }, ptr addrspace(6) %4, i32 0, i32 0 + %6 = load <2 x float>, ptr addrspace(6) %5, align 8 + %7 = call ptr addrspace(1) @_Z14clspv.resource.4(i32 0, i32 4, i32 0, i32 4, i32 4, i32 0, { [0 x <4 x float>] } zeroinitializer) + %8 = getelementptr { [0 x <4 x float>] }, ptr addrspace(1) %7, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x float> @_Z11read_imagef30ocl_image2d_ro_t.float.sampled11ocl_samplerDv2_f(target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0) %0, target("spirv.Sampler") %1, <2 x float> %6) + %9 = load <4 x float>, ptr addrspace(2) %3, align 16 + %add = fadd <4 x float> %call, %9 + store <4 x float> %add, ptr addrspace(1) %8, align 16 + ret void +} + +declare target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, target("spirv.Image", float, 1, 0, 0, 0, 1, 0, 0, 0)) + +declare target("spirv.Sampler") @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, target("spirv.Sampler")) + +declare ptr addrspace(2) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { [4096 x <4 x float>] }) + +declare ptr addrspace(6) @_Z14clspv.resource.3(i32, i32, i32, i32, i32, i32, { <2 x float> }) + +declare ptr addrspace(1) @_Z14clspv.resource.4(i32, i32, i32, i32, i32, i32, { [0 x <4 x float>] }) + +!8 = !{i32 1} + diff --git a/test/UBO/constant_wrapping.cl b/test/UBO/constant_wrapping.cl deleted file mode 100644 index 696569c60..000000000 --- a/test/UBO/constant_wrapping.cl +++ /dev/null @@ -1,56 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -typedef struct { - int x; - int y __attribute((aligned(16))); -} inner; - -typedef struct { - inner i[2]; -} outer; - -__kernel void foo(__global inner* data, __constant outer* c) { - unsigned gid = get_global_id(0); - data[gid].x = c[gid].i[0].x; - data[gid].y = c[gid].i[0].y; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[inner:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[inner]] 1 Offset 4 -// CHECK-DAG: OpMemberDecorate [[inner]] 2 Offset 16 -// CHECK-DAG: OpMemberDecorate [[inner]] 3 Offset 20 -// CHECK-DAG: OpDecorate [[inner_runtime:%[0-9a-zA-Z_]+]] ArrayStride 32 -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 64 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[inner]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] -// CHECK-DAG: [[inner_runtime]] = OpTypeRuntimeArray [[inner]] -// CHECK-DAG: [[data_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[inner_runtime]] -// CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[data_struct]] -// CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 -// CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[inner]] [[two]] -// CHECK-DAG: [[outer:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 -// CHECK-DAG: [[ubo_array:%[0-9a-zA-Z_]+]] = OpTypeArray [[outer]] [[int_1024]] -// CHECK-DAG: [[block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[block]] -// CHECK-DAG: [[c_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] -// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c]] = OpVariable [[c_ptr]] Uniform -// CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_ele_ptr]] [[c]] [[zero]] {{.*}} [[zero]] [[zero]] [[zero]] -// CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_gep]] -// CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] {{.*}} -// CHECK: OpStore [[data_gep]] [[c_load]] diff --git a/test/UBO/constant_wrapping.ll b/test/UBO/constant_wrapping.ll new file mode 100644 index 000000000..8f20b86fc --- /dev/null +++ b/test/UBO/constant_wrapping.ll @@ -0,0 +1,74 @@ +; RUN: clspv-opt -constant-args-ubo -int8=0 %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[inner:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[inner]] 1 Offset 4 +; CHECK-DAG: OpMemberDecorate [[inner]] 2 Offset 16 +; CHECK-DAG: OpMemberDecorate [[inner]] 3 Offset 20 +; CHECK-DAG: OpDecorate [[inner_runtime:%[0-9a-zA-Z_]+]] ArrayStride 32 +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 64 +; CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[inner]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] +; CHECK-DAG: [[inner_runtime]] = OpTypeRuntimeArray [[inner]] +; CHECK-DAG: [[data_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[inner_runtime]] +; CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[data_struct]] +; CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 +; CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[inner]] [[two]] +; CHECK-DAG: [[outer:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 +; CHECK-DAG: [[ubo_array:%[0-9a-zA-Z_]+]] = OpTypeArray [[outer]] [[int_1024]] +; CHECK-DAG: [[block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[block]] +; CHECK-DAG: [[c_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] +; CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] +; CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[c]] = OpVariable [[c_ptr]] Uniform +; CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_ele_ptr]] [[c]] [[zero]] {{.*}} [[zero]] [[zero]] [[zero]] +; CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_gep]] +; CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] {{.*}} +; CHECK: OpStore [[data_gep]] [[c_load]] + +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" + +%struct.inner = type { i32, [12 x i8], i32, [12 x i8] } +%struct.outer = type { [2 x %struct.inner] } + +@__spirv_GlobalInvocationId = local_unnamed_addr addrspace(5) global <3 x i32> zeroinitializer +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !8 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.inner] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [1024 x %struct.outer] } zeroinitializer) + %2 = getelementptr <3 x i32>, ptr addrspace(5) @__spirv_GlobalInvocationId, i32 0, i32 0 + %3 = load i32, ptr addrspace(5) %2, align 16 + %4 = getelementptr { [1024 x %struct.outer] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 0 + %5 = load i32, ptr addrspace(2) %4, align 16 + %6 = getelementptr { [0 x %struct.inner] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 0 + store i32 %5, ptr addrspace(1) %6, align 16 + %7 = getelementptr { [1024 x %struct.outer] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 2 + %8 = load i32, ptr addrspace(2) %7, align 16 + %9 = getelementptr { [0 x %struct.inner] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 2 + store i32 %8, ptr addrspace(1) %9, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.inner] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [1024 x %struct.outer] }) + +!8 = !{i32 1} + diff --git a/test/UBO/copy.cl b/test/UBO/copy.cl deleted file mode 100644 index 45ad7cc92..000000000 --- a/test/UBO/copy.cl +++ /dev/null @@ -1,28 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -__kernel void foo(__global int4* data, __constant int4* c) { - *data = *c; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpDecorate [[var:%[0-9a-zA-Z_]+]] NonWritable -// CHECK-DAG: OpDecorate [[var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[var]] Binding 1 -// CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK: [[int4:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 4 -// CHECK: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[int4]] [[int_4096]] -// CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK: [[ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct]] -// CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK: [[ptr_int4:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int4]] -// CHECK: [[var]] = OpVariable [[ptr]] Uniform -// CHECK: [[gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_int4]] [[var]] [[zero]] [[zero]] -// CHECK: OpLoad [[int4]] [[gep]] diff --git a/test/UBO/copy.ll b/test/UBO/copy.ll new file mode 100644 index 000000000..384fb9fab --- /dev/null +++ b/test/UBO/copy.ll @@ -0,0 +1,47 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpDecorate [[var:%[0-9a-zA-Z_]+]] NonWritable +; CHECK-DAG: OpDecorate [[var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[var]] Binding 1 +; CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK: [[int4:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 4 +; CHECK: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[int4]] [[int_4096]] +; CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK: [[ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct]] +; CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK: [[ptr_int4:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int4]] +; CHECK: [[var]] = OpVariable [[ptr]] Uniform +; CHECK: [[gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_int4]] [[var]] [[zero]] [[zero]] +; CHECK: OpLoad [[int4]] [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = load <4 x i32>, ptr addrspace(2) %3, align 16 + store <4 x i32> %4, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!14 = !{i32 2} + diff --git a/test/UBO/copy_nested.cl b/test/UBO/copy_nested.cl deleted file mode 100644 index 01cf9cdd2..000000000 --- a/test/UBO/copy_nested.cl +++ /dev/null @@ -1,40 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -typedef struct inner { - float4 x; - float4 y; -} inner; - -typedef struct outer { - inner x; -} outer; - -__kernel void foo(__global outer* data, __constant outer* c) { - data->x.x = c->x.x; - data->x.y = c->x.y; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpDecorate [[var:%[0-9a-zA-Z_]+]] NonWritable -// CHECK-DAG: OpDecorate [[var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[var]] Binding 1 -// CHECK: OpDecorate [[array:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: [[float:%[0-9a-zA-Z_]+]] = OpTypeFloat 32 -// CHECK-DAG: [[float4:%[0-9a-zA-Z_]+]] = OpTypeVector [[float]] 4 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK-DAG: [[array]] = OpTypeArray [[float4]] [[int_4096]] -// CHECK-DAG: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK-DAG: [[ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct]] -// CHECK-DAG: [[ptr_float4:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[float4]] -// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK: [[var]] = OpVariable [[ptr]] Uniform -// CHECK: [[gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_float4]] [[var]] [[zero]] [[zero]] -// CHECK: OpLoad [[float4]] [[gep]] diff --git a/test/UBO/copy_nested.ll b/test/UBO/copy_nested.ll new file mode 100644 index 000000000..526b4e7e0 --- /dev/null +++ b/test/UBO/copy_nested.ll @@ -0,0 +1,53 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpDecorate [[var:%[0-9a-zA-Z_]+]] NonWritable +; CHECK-DAG: OpDecorate [[var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[var]] Binding 1 +; CHECK: OpDecorate [[array:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: [[float:%[0-9a-zA-Z_]+]] = OpTypeFloat 32 +; CHECK-DAG: [[float4:%[0-9a-zA-Z_]+]] = OpTypeVector [[float]] 4 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK-DAG: [[array]] = OpTypeArray [[float4]] [[int_4096]] +; CHECK-DAG: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK-DAG: [[ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct]] +; CHECK-DAG: [[ptr_float4:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[float4]] +; CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK: [[var]] = OpVariable [[ptr]] Uniform +; CHECK: [[gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_float4]] [[var]] [[zero]] [[zero]] +; CHECK: OpLoad [[float4]] [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x float>] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x float>] } zeroinitializer) + %2 = getelementptr { [4096 x <4 x float>] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0 + %3 = load <4 x float>, ptr addrspace(2) %2, align 16 + %4 = getelementptr { [0 x <4 x float>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + store <4 x float> %3, ptr addrspace(1) %4, align 16 + %5 = getelementptr { [4096 x <4 x float>] }, ptr addrspace(2) %1, i32 0, i32 0, i32 1 + %6 = load <4 x float>, ptr addrspace(2) %5, align 16 + %7 = getelementptr { [0 x <4 x float>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 1 + store <4 x float> %6, ptr addrspace(1) %7, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x float>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x float>] }) + +!14 = !{i32 2} + diff --git a/test/UBO/extra_arg.cl b/test/UBO/extra_arg.cl deleted file mode 100644 index 20af3c318..000000000 --- a/test/UBO/extra_arg.cl +++ /dev/null @@ -1,36 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -arch=spir -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm --check-prefixes=CHECK,CHECK-32 -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -arch=spir64 -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm --check-prefixes=CHECK,CHECK-64 -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -__attribute__((noinline)) -int4 bar(constant int4* in) { return in[0]; } - -kernel void k1(global int4* out, constant int4* in) { - constant int4* x = in + in[0].x; - *out = bar(x); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK-64: [[long:%[a-zA-Z0-9_]+]] = OpTypeInt 64 0 -// CHECK: [[var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK: [[k1]] = OpFunction -// CHECK: [[ex:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] {{.*}} 0 -// CHECK-64: [[ex_long:%[a-zA-Z0-9_]+]] = OpSConvert [[long]] [[ex]] -// CHECK-64: OpFunctionCall {{.*}} [[bar:%[a-zA-Z0-9_]+]] [[ex_long]] -// CHECK-32: OpFunctionCall {{.*}} [[bar:%[a-zA-Z0-9_]+]] [[ex]] -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[bar]] = OpFunction -// CHECK-64-NEXT: [[param:%[a-zA-Z0-9_]+]] = OpFunctionParameter [[long]] -// CHECK-32-NEXT: [[param:%[a-zA-Z0-9_]+]] = OpFunctionParameter [[int]] -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[param]] -// CHECK: OpLoad {{.*}} [[gep]] -// CHECK: OpFunctionEnd -// CHECK-NOT: OpFunction diff --git a/test/UBO/extra_arg.ll b/test/UBO/extra_arg.ll new file mode 100644 index 000000000..db151cfd4 --- /dev/null +++ b/test/UBO/extra_arg.ll @@ -0,0 +1,42 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} @k1 +; CHECK: [[val:%[a-zA-Z0-9_.]+]] = extractelement <4 x i32> {{.*}}, i64 0 +; CHECK: call {{.*}} [[bar:@bar[a-zA-Z0-9_.]+]](i32 [[val]]) +; CHECK: define {{.*}} [[bar]](i32 [[p:%[a-zA-Z0-9_.]+]]) +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 [[p]] +; CHECK: load <4 x i32>, ptr addrspace(2) [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %in) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %in, align 16 + ret <4 x i32> %0 +} + +define dso_local spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !16 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = load <4 x i32>, ptr addrspace(2) %3, align 16 + %5 = extractelement <4 x i32> %4, i64 0 + %6 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 %5 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %6) + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!16 = !{i32 2} + diff --git a/test/UBO/extra_args.cl b/test/UBO/extra_args.cl deleted file mode 100644 index 83dd0c457..000000000 --- a/test/UBO/extra_args.cl +++ /dev/null @@ -1,45 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -arch=spir -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm --check-prefixes=CHECK,CHECK-32 -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -arch=spir64 -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm --check-prefixes=CHECK,CHECK-64 -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -__attribute__((noinline)) -int4 bar(constant int4* in1, constant int4* in2) { - return in1[0] + in2[0]; -} - -kernel void k1(global int4* out, constant int4* in) { - constant int4* x = in + in[0].x; - constant int4* y = in + in[1].y; - *out = bar(x, y); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK-64: [[long:%[a-zA-Z0-9_]+]] = OpTypeInt 64 0 -// CHECK: [[var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK: [[k1]] = OpFunction -// CHECK: [[ex0:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] {{.*}} 0 -// CHECK-64: [[ex0_long:%[a-zA-Z0-9_]+]] = OpSConvert [[long]] [[ex0]] -// CHECK: [[ex1:%[a-zA-Z0-9_]+]] = OpCompositeExtract [[int]] {{.*}} 1 -// CHECK-64: [[ex1_long:%[a-zA-Z0-9_]+]] = OpSConvert [[long]] [[ex1]] -// CHECK-64: OpFunctionCall {{.*}} [[bar:%[a-zA-Z0-9_]+]] [[ex0_long]] [[ex1_long]] -// CHECK-32: OpFunctionCall {{.*}} [[bar:%[a-zA-Z0-9_]+]] [[ex0]] [[ex1]] -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[bar]] = OpFunction -// CHECK-64-NEXT: [[param0:%[a-zA-Z0-9_]+]] = OpFunctionParameter [[long]] -// CHECK-64-NEXT: [[param1:%[a-zA-Z0-9_]+]] = OpFunctionParameter [[long]] -// CHECK-32-NEXT: [[param0:%[a-zA-Z0-9_]+]] = OpFunctionParameter [[int]] -// CHECK-32-NEXT: [[param1:%[a-zA-Z0-9_]+]] = OpFunctionParameter [[int]] -// CHECK: [[gep0:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[param0]] -// CHECK: [[gep1:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[param1]] -// CHECK: OpLoad {{.*}} [[gep0]] -// CHECK: OpLoad {{.*}} [[gep1]] -// CHECK: OpFunctionEnd -// CHECK-NOT: OpFunction diff --git a/test/UBO/extra_args.ll b/test/UBO/extra_args.ll new file mode 100644 index 000000000..0a0de1bdc --- /dev/null +++ b/test/UBO/extra_args.ll @@ -0,0 +1,50 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} @k1 +; CHECK-DAG: [[ex0:%[a-zA-Z0-9_.]+]] = extractelement <4 x i32> {{.*}}, i64 0 +; CHECK-DAG: [[ex1:%[a-zA-Z0-9_.]+]] = extractelement <4 x i32> {{.*}}, i64 1 +; CHECK: call {{.*}} [[bar:@bar[a-zA-Z0-9_.]+]](i32 [[ex0]], i32 [[ex1]]) +; CHECK: define {{.*}} [[bar]](i32 [[v0:%[a-zA-Z0-9_.]+]], i32 [[v1:%[a-zA-Z0-9_.]+]]) +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 [[v0]] +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource.1 +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 [[v1]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %in1, ptr addrspace(2) nocapture readonly %in2) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %in1, align 16 + %1 = load <4 x i32>, ptr addrspace(2) %in2, align 16 + %add = add <4 x i32> %1, %0 + ret <4 x i32> %add +} + +define dso_local spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !16 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = load <4 x i32>, ptr addrspace(2) %3, align 16 + %5 = extractelement <4 x i32> %4, i64 0 + %6 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 %5 + %7 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 1 + %8 = load <4 x i32>, ptr addrspace(2) %7, align 16 + %9 = extractelement <4 x i32> %8, i64 1 + %10 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 %9 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %6, ptr addrspace(2) %10) #2 + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!16 = !{i32 2} + diff --git a/test/UBO/global_wrapping.cl b/test/UBO/global_wrapping.cl deleted file mode 100644 index 444ecefdb..000000000 --- a/test/UBO/global_wrapping.cl +++ /dev/null @@ -1,58 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -typedef struct { - int x; - int y __attribute((aligned(16))); -} inner; - -typedef struct { - inner i[2]; -} outer; - -__kernel void foo(__global outer* data, __constant inner* c) { - unsigned gid = get_global_id(0); - data[gid].i[0].x = c[gid].x; - data[gid].i[0].y = c[gid].y; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[inner:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[inner]] 1 Offset 4 -// CHECK-DAG: OpMemberDecorate [[inner]] 2 Offset 16 -// CHECK-DAG: OpMemberDecorate [[inner]] 3 Offset 20 -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 64 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[array:%[0-9a-zA-Z_]+]] ArrayStride 32 -// CHECK-DAG: OpDecorate [[inner_array:%[0-9a-zA-Z_]+]] ArrayStride 32 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[inner]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] -// CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 -// CHECK-DAG: [[array]] = OpTypeArray [[inner]] [[two]] -// CHECK-DAG: [[outer:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[outer]] -// CHECK-DAG: [[block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[block]] -// CHECK-DAG: [[int_2048:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2048 -// CHECK-DAG: [[inner_array]] = OpTypeArray [[inner]] [[int_2048]] -// CHECK-DAG: [[c_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[inner_array]] -// CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[c_struct]] -// CHECK-DAG: [[c_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] -// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c]] = OpVariable [[c_ptr]] Uniform -// CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_ele_ptr]] [[c]] [[zero]] {{.*}} [[zero]] -// CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_gep]] -// CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] {{.*}} [[zero]] [[zero]] -// CHECK: OpStore [[data_gep]] [[c_load]] - diff --git a/test/UBO/global_wrapping.ll b/test/UBO/global_wrapping.ll new file mode 100644 index 000000000..31cfe8965 --- /dev/null +++ b/test/UBO/global_wrapping.ll @@ -0,0 +1,75 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t -producer-out-file %t.spv -int8=0 --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[inner:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[inner]] 1 Offset 4 +; CHECK-DAG: OpMemberDecorate [[inner]] 2 Offset 16 +; CHECK-DAG: OpMemberDecorate [[inner]] 3 Offset 20 +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 64 +; CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[array:%[0-9a-zA-Z_]+]] ArrayStride 32 +; CHECK-DAG: OpDecorate [[inner_array:%[0-9a-zA-Z_]+]] ArrayStride 32 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[inner]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] +; CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 +; CHECK-DAG: [[array]] = OpTypeArray [[inner]] [[two]] +; CHECK-DAG: [[outer:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[outer]] +; CHECK-DAG: [[block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] +; CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[block]] +; CHECK-DAG: [[int_2048:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2048 +; CHECK-DAG: [[inner_array]] = OpTypeArray [[inner]] [[int_2048]] +; CHECK-DAG: [[c_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[inner_array]] +; CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[c_struct]] +; CHECK-DAG: [[c_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] +; CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] +; CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[c]] = OpVariable [[c_ptr]] Uniform +; CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_ele_ptr]] [[c]] [[zero]] {{.*}} [[zero]] +; CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_gep]] +; CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] {{.*}} [[zero]] [[zero]] +; CHECK: OpStore [[data_gep]] [[c_load]] + +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" + +%struct.outer = type { [2 x %struct.inner] } +%struct.inner = type { i32, [12 x i8], i32, [12 x i8] } + +@__spirv_GlobalInvocationId = local_unnamed_addr addrspace(5) global <3 x i32> zeroinitializer +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !13 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.outer] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [2048 x %struct.inner] } zeroinitializer) + %2 = getelementptr <3 x i32>, ptr addrspace(5) @__spirv_GlobalInvocationId, i32 0, i32 0 + %3 = load i32, ptr addrspace(5) %2, align 16 + %4 = getelementptr { [2048 x %struct.inner] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 0 + %5 = load i32, ptr addrspace(2) %4, align 16 + %6 = getelementptr { [0 x %struct.outer] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 0 + store i32 %5, ptr addrspace(1) %6, align 16 + %7 = getelementptr { [2048 x %struct.inner] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 2 + %8 = load i32, ptr addrspace(2) %7, align 16 + %9 = getelementptr { [0 x %struct.outer] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 2 + store i32 %8, ptr addrspace(1) %9, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.outer] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [2048 x %struct.inner] }) + +!13 = !{i32 2} + diff --git a/test/UBO/large_padding.cl b/test/UBO/large_padding.cl deleted file mode 100644 index db00b9002..000000000 --- a/test/UBO/large_padding.cl +++ /dev/null @@ -1,41 +0,0 @@ -// RUN: clspv %target -int8 -constant-args-ubo -inline-entry-points %s -o %t.spv -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: clspv-reflection %t.spv -o %t.map -// RUN: FileCheck -check-prefix=MAP %s < %t.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// TODO(#1292) -// XFAIL: * - -// Prior to #279 this would produce a [16 x i8] padding array. -typedef struct { - int4 a; - int4 b __attribute__((aligned(32))); - int4 c; -} S; - -kernel void foo(global S* out, constant S* in) { - out->c = in->c; -} - -// MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 16 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 32 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 48 -// CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[in]] NonWritable -// CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 64 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 -// CHECK-DAG: [[int4:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 4 -// CHECK: [[s]] = OpTypeStruct [[int4]] [[char]] [[int4]] [[int4]] -// CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 -// CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_1024]] -// CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] -// CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform diff --git a/test/UBO/large_padding.ll b/test/UBO/large_padding.ll new file mode 100644 index 000000000..b3d751aa1 --- /dev/null +++ b/test/UBO/large_padding.ll @@ -0,0 +1,54 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 16 +; CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 32 +; CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 48 +; CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[in]] NonWritable +; CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 64 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 +; CHECK-DAG: [[int4:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 4 +; CHECK: [[s]] = OpTypeStruct [[int4]] [[char]] [[int4]] [[int4]] +; CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 +; CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_1024]] +; CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] +; CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform + +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" + +%struct.S = type { <4 x i32>, [16 x i8], <4 x i32>, <4 x i32> } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 32 %out, ptr addrspace(2) nocapture readonly align 32 %in) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.S] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [1024 x %struct.S] } zeroinitializer) + %2 = getelementptr { [1024 x %struct.S] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0, i32 3 + %3 = load <4 x i32>, ptr addrspace(2) %2, align 16 + %4 = getelementptr { [0 x %struct.S] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0, i32 3 + store <4 x i32> %3, ptr addrspace(1) %4, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.S] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [1024 x %struct.S] }) + +declare ptr addrspace(9) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { { i32 } }) + +!14 = !{i32 2} + diff --git a/test/UBO/large_padding_std430.cl b/test/UBO/large_padding_std430.cl deleted file mode 100644 index f9af9561d..000000000 --- a/test/UBO/large_padding_std430.cl +++ /dev/null @@ -1,45 +0,0 @@ -// RUN: clspv %target -int8 -constant-args-ubo -inline-entry-points -std430-ubo-layout %s -o %t.spv -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: clspv-reflection -d %t.spv -o %t.map -// RUN: FileCheck -check-prefix=MAP %s < %t.map - -// TODO(#1292) -// XFAIL: * - -// With std430 layouts in UBO, the padding array ([16 x i8]) can be generated -// with an ArrayStride of 1. -typedef struct { - int4 a; - int4 b __attribute__((aligned(32))); - int4 c; -} S; - -kernel void foo(global S* out, constant S* in) { - out->c = in->c; -} - -// MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 16 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 32 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 48 -// CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[in]] NonWritable -// CHECK: OpDecorate [[char_array:%[0-9a-zA-Z_]+]] ArrayStride 1 -// CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 64 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 -// CHECK-DAG: [[int4:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 4 -// CHECK-DAG: [[int_16:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 16 -// CHECK-DAG: [[char_array]] = OpTypeArray [[char]] [[int_16]] -// CHECK: [[s]] = OpTypeStruct [[int4]] [[char_array]] [[int4]] [[int4]] -// CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 -// CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_1024]] -// CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] -// CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform - diff --git a/test/UBO/large_padding_std430.ll b/test/UBO/large_padding_std430.ll new file mode 100644 index 000000000..f112e386c --- /dev/null +++ b/test/UBO/large_padding_std430.ll @@ -0,0 +1,59 @@ +; RUN: clspv-opt -constant-args-ubo -std430-ubo-layout %s -o %t -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv --uniform-buffer-standard-layout +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,out,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,in,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; With std430 layouts in UBO, the padding array ([16 x i8]) can be generated +; with an ArrayStride of 1. +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 16 +; CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 32 +; CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 48 +; CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[in]] NonWritable +; CHECK: OpDecorate [[char_array:%[0-9a-zA-Z_]+]] ArrayStride 1 +; CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 64 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 +; CHECK-DAG: [[int4:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 4 +; CHECK-DAG: [[int_16:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 16 +; CHECK-DAG: [[char_array]] = OpTypeArray [[char]] [[int_16]] +; CHECK: [[s]] = OpTypeStruct [[int4]] [[char_array]] [[int4]] [[int4]] +; CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 +; CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_1024]] +; CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] +; CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform + +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" + +%struct.S = type { <4 x i32>, [16 x i8], <4 x i32>, <4 x i32> } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 32 %out, ptr addrspace(2) nocapture readonly align 32 %in) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.S] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [1024 x %struct.S] } zeroinitializer) + %5 = getelementptr { [1024 x %struct.S] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0, i32 3 + %6 = load <4 x i32>, ptr addrspace(2) %5, align 16 + %7 = getelementptr { [0 x %struct.S] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0, i32 3 + store <4 x i32> %6, ptr addrspace(1) %7, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.S] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [1024 x %struct.S] }) + +declare ptr addrspace(9) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { { i32 } }) + +!14 = !{i32 2} + diff --git a/test/UBO/long_specialization_chain.cl b/test/UBO/long_specialization_chain.cl deleted file mode 100644 index 90bc2d8f7..000000000 --- a/test/UBO/long_specialization_chain.cl +++ /dev/null @@ -1,74 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -// TODO(#1292) -// XFAIL: * - -__attribute__((noinline)) -int4 c(constant int4* data) { return data[0]; } - -__attribute__((noinline)) -int4 b(constant int4* data) { return c(data); } - -__attribute__((noinline)) -int4 a(constant int4* data) { return b(data); } - -kernel void k1(global int4* out, constant int4* in) { - *out = a(in); -} - -kernel void k2(global int4* out, constant int4* in) { - *out = a(in + 1); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: OpEntryPoint GLCompute [[k2:%[a-zA-Z0-9_]+]] -// CHECK-DAG: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[var:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK-DAG: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[int1:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 1 -// CHECK: [[k1]] = OpFunction -// CHECK: [[call:%[a-zA-Z0-9_]+]] = OpFunctionCall {{.*}} [[a1:%[a-zA-Z0-9_]+]] -// CHECK-NEXT: OpStore {{.*}} [[call]] -// CHECK-NEXT: OpReturn -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[k2]] = OpFunction -// CHECK: [[call:%[a-zA-Z0-9_]+]] = OpFunctionCall {{.*}} [[a2:%[a-zA-Z0-9_]+]] -// CHECK-NEXT: OpStore {{.*}} [[call]] -// CHECK-NEXT: OpReturn -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[a2]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[call:%[a-zA-Z0-9_]+]] = OpFunctionCall {{.*}} [[b2:%[a-zA-Z0-9_]+]] -// CHECK: ReturnValue [[call]] -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[a1]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[call:%[a-zA-Z0-9_]+]] = OpFunctionCall {{.*}} [[b1:%[a-zA-Z0-9_]+]] -// CHECK: ReturnValue [[call]] -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[b1]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[call:%[a-zA-Z0-9_]+]] = OpFunctionCall {{.*}} [[c1:%[a-zA-Z0-9_]+]] -// CHECK: ReturnValue [[call]] -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[b2]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[call:%[a-zA-Z0-9_]+]] = OpFunctionCall {{.*}} [[c2:%[a-zA-Z0-9_]+]] -// CHECK: ReturnValue [[call]] -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[c2]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int1]] -// CHECK: [[ld:%[a-zA-Z0-9_]+]] = OpLoad {{.*}} [[gep]] -// CHECK: ReturnValue [[ld]] -// CHECK-NEXT: OpFunctionEnd -// CHECK-NEXT: [[c1]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var]] [[int0]] [[int0]] -// CHECK: [[ld:%[a-zA-Z0-9_]+]] = OpLoad {{.*}} [[gep]] -// CHECK: ReturnValue [[ld]] -// CHECK-NEXT: OpFunctionEnd -// CHECK-NOT: OpFunction diff --git a/test/UBO/long_specialization_chain.ll b/test/UBO/long_specialization_chain.ll new file mode 100644 index 000000000..93190adb0 --- /dev/null +++ b/test/UBO/long_specialization_chain.ll @@ -0,0 +1,76 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} @k1 +; CHECK: call {{.*}} [[a1:@a[a-zA-Z0-9_.]+]] +; CHECK: define {{.*}} @k2 +; CHECK: call {{.*}} [[a2:@a[a-zA-Z0-9_.]+]] + +; CHECK: define {{.*}} [[a2]] +; CHECK: call {{.*}} [[b2:@b[a-zA-Z0-9_.]+]] +; CHECK: define {{.*}} [[a1]] +; CHECK: call {{.*}} [[b1:@b[a-zA-Z0-9_.]+]] + +; CHECK: define {{.*}} [[b1]] +; CHECK: call {{.*}} [[c1:@c[a-zA-Z0-9_.]+]] +; CHECK: define {{.*}} [[b2]] +; CHECK: call {{.*}} [[c2:@c[a-zA-Z0-9_.]+]] + +; CHECK: define {{.*}} [[c2]]() +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 1 + +; CHECK: define {{.*}} [[c1]]() +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 0 + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @c(ptr addrspace(2) nocapture readonly %data) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %data, align 16 + ret <4 x i32> %0 +} + +define dso_local spir_func <4 x i32> @b(ptr addrspace(2) nocapture readonly %data) { +entry: + %call = tail call spir_func <4 x i32> @c(ptr addrspace(2) %data) + ret <4 x i32> %call +} + +define dso_local spir_func <4 x i32> @a(ptr addrspace(2) nocapture readonly %data) { +entry: + %call = tail call spir_func <4 x i32> @b(ptr addrspace(2) %data) + ret <4 x i32> %call +} + +define dso_local spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !19 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @a(ptr addrspace(2) %3) + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +define spir_kernel void @k2(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !19 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %6 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 1 + %call.i = tail call spir_func <4 x i32> @a(ptr addrspace(2) %6) + store <4 x i32> %call.i, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!19 = !{i32 2} diff --git a/test/UBO/max_ubo_size.cl b/test/UBO/max_ubo_size.cl deleted file mode 100644 index 34bb936d2..000000000 --- a/test/UBO/max_ubo_size.cl +++ /dev/null @@ -1,45 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -max-ubo-size=64 -int8=0 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// Checking that -max-ubo-size affects the number of elements in the UBO array. -// Struct alloca size is 32, so expect 2 elements with max size of 64. - -typedef struct { - int x; - int y __attribute((aligned(16))); -} s; - -__kernel void foo(__global s* data, __constant s* c) { - unsigned gid = get_global_id(0); - data[gid].x = c[gid].x; - data[gid].y = c[gid].y; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 4 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 16 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 20 -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 32 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 -// CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK: [[s]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] -// CHECK: [[runtime]] = OpTypeRuntimeArray [[s]] -// CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] -// CHECK: [[int_2:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 -// CHECK: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[s]] [[int_2]] -// CHECK: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c]] = OpVariable [[c_ptr]] Uniform - diff --git a/test/UBO/max_ubo_size.ll b/test/UBO/max_ubo_size.ll new file mode 100644 index 000000000..e06844e7d --- /dev/null +++ b/test/UBO/max_ubo_size.ll @@ -0,0 +1,34 @@ +; RUN: clspv-opt -constant-args-ubo -max-ubo-size=64 %s -o %t.ll --passes=allocate-descriptors +; RUN: FileCheck %s < %t.ll + +; Checking that -max-ubo-size affects the number of elements in the UBO array. +; Struct alloca size is 32, so expect 2 elements with max size of 64. +; CHECK: [[s:%[a-zA-Z0-9_.]+]] = type { i32, [12 x i8], i32, [12 x i8] } +; CHECK: call ptr addrspace(2) @_Z14clspv.resource +; CHECK-SAME: { [2 x [[s]]] } zeroinitializer + +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" + +%struct.s = type { i32, [12 x i8], i32, [12 x i8] } + +@__spirv_GlobalInvocationId = local_unnamed_addr addrspace(5) global <3 x i32> zeroinitializer +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !12 { +entry: + %0 = getelementptr <3 x i32>, ptr addrspace(5) @__spirv_GlobalInvocationId, i32 0, i32 0 + %1 = load i32, ptr addrspace(5) %0, align 16 + %2 = getelementptr inbounds %struct.s, ptr addrspace(2) %c, i32 %1, i32 0 + %3 = load i32, ptr addrspace(2) %2, align 16 + %4 = getelementptr inbounds %struct.s, ptr addrspace(1) %data, i32 %1, i32 0 + store i32 %3, ptr addrspace(1) %4, align 16 + %5 = getelementptr inbounds %struct.s, ptr addrspace(2) %c, i32 %1, i32 2 + %6 = load i32, ptr addrspace(2) %5, align 16 + %7 = getelementptr inbounds %struct.s, ptr addrspace(1) %data, i32 %1, i32 2 + store i32 %6, ptr addrspace(1) %7, align 16 + ret void +} + +!12 = !{i32 2} + diff --git a/test/UBO/mixed_inlining.cl b/test/UBO/mixed_inlining.cl deleted file mode 100644 index 8ae836cf5..000000000 --- a/test/UBO/mixed_inlining.cl +++ /dev/null @@ -1,35 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm - -// This test is not validated because it uses a selection between Uniform -// pointers, which is disallowed by SPIR-V. -__attribute__((noinline)) -int4 bar(constant int4* data) { return data[0]; } - -kernel void k1(global int4* out, constant int4* in1, constant int4* in2, int a) { - constant int4* x = (a == 0) ? in1 : in2; - // This call requires inlining. - *out = bar(x); -} - -kernel void k2(global int4* out, constant int4* in) { - // This call is specialized. - *out = bar(in); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: OpEntryPoint GLCompute [[k2:%[a-zA-Z0-9_]+]] -// CHECK: [[k1]] = OpFunction -// CHECK-NOT: OpFunctionCall -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[k2]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar:%[a-zA-Z0-9_]+]] -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[bar]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep:%[a-zA-Z0-9_]+]] = OpAccessChain -// CHECK: OpLoad {{.*}} [[gep]] -// CHECK: OpFunctionEnd -// CHECK-NOT: OpFunction - diff --git a/test/UBO/mixed_inlining.ll b/test/UBO/mixed_inlining.ll new file mode 100644 index 000000000..478725abe --- /dev/null +++ b/test/UBO/mixed_inlining.ll @@ -0,0 +1,75 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; Note: this test would produce invalid SPIR-V. + +; The call to bar in k1 requires inlining +; CHECK: define {{.*}} @k1 +; CHECK: [[sel:%[a-zA-Z0-9_.]+]] = select i1 %{{[^ ]+}}, ptr addrspace(2) +; CHECK: load <4 x i32>, ptr addrspace(2) [[sel]] + +; The call to bar in k2 is specialized +; CHECK: define {{.*}} @k2 +; CHECK: call {{.*}} [[bar2:@bar[a-zA-Z0-9_.]+]]() + +; CHECK: define {{.*}} [[bar2]] +; CHECK: [[res:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK: [[gep:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res]], i32 0, i32 0, i32 0 +; CHECK: load <4 x i32>, ptr addrspace(2) [[gep]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %data) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %data, align 16 + ret <4 x i32> %0 +} + +define spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in1, ptr addrspace(2) nocapture readonly align 16 %in2, { i32 } %podargs) !clspv.pod_args_impl !18 !kernel_arg_map !19 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = call ptr addrspace(2) @_Z14clspv.resource.2(i32 0, i32 2, i32 1, i32 2, i32 2, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %5 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %4, i32 0, i32 0, i32 0 + %6 = call ptr addrspace(9) @_Z14clspv.resource.3(i32 -1, i32 3, i32 5, i32 3, i32 3, i32 0, { { i32 } } zeroinitializer) + %7 = getelementptr { { i32 } }, ptr addrspace(9) %6, i32 0, i32 0 + %8 = load { i32 }, ptr addrspace(9) %7, align 4 + %a = extractvalue { i32 } %8, 0 + %cmp.i = icmp eq i32 %a, 0 + %in1.in2 = select i1 %cmp.i, ptr addrspace(2) %3, ptr addrspace(2) %5 + %call.i = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %in1.in2) #2 + store <4 x i32> %call.i, ptr addrspace(1) %1, align 16 + ret void +} + +define dso_local spir_kernel void @k2(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in) !clspv.pod_args_impl !18 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %3) #2 + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +declare ptr addrspace(9) @_Z14clspv.resource.3(i32, i32, i32, i32, i32, i32, { { i32 } }) + +!18 = !{i32 2} +!19 = !{!20, !21, !22, !23} +!20 = !{!"out", i32 0, i32 0, i32 0, i32 0, !"buffer"} +!21 = !{!"in1", i32 1, i32 1, i32 0, i32 0, !"buffer_ubo"} +!22 = !{!"in2", i32 2, i32 2, i32 0, i32 0, !"buffer_ubo"} +!23 = !{!"a", i32 3, i32 3, i32 0, i32 4, !"pod_pushconstant"} + diff --git a/test/UBO/multiple_ubo_args.cl b/test/UBO/multiple_ubo_args.cl deleted file mode 100644 index 1575496b3..000000000 --- a/test/UBO/multiple_ubo_args.cl +++ /dev/null @@ -1,47 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: spirv-val %t.spv --target-env vulkan1.0 - -__attribute__((noinline)) -int4 bar(constant int4* in1, constant int4* in2) { - return in1[0] + in2[0]; -} - -kernel void k1(global int4* out, constant int4* in1, constant int4* in2) { - *out = bar(in1, in2); -} - -kernel void k2(global int4* out, constant int4* in1, constant int4* in2) { - *out = bar(in2, in1); -} - -// CHECK: OpEntryPoint GLCompute [[k1:%[a-zA-Z0-9_]+]] -// CHECK: OpEntryPoint GLCompute [[k2:%[a-zA-Z0-9_]+]] -// CHECK: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 -// CHECK: [[int0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK: [[var1:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK: [[var2:%[a-zA-Z0-9_]+]] = OpVariable {{.*}} Uniform -// CHECK: [[k1]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar1:%[a-zA-Z0-9_]+]] -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[k2]] = OpFunction -// CHECK: OpFunctionCall {{.*}} [[bar2:%[a-zA-Z0-9_]+]] -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[bar2]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep1:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var2]] [[int0]] [[int0]] -// CHECK: [[gep2:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var1]] [[int0]] [[int0]] -// CHECK: [[ld1:%[a-zA-Z0-9_]+]] = OpLoad {{.*}} [[gep1]] -// CHECK: [[ld2:%[a-zA-Z0-9_]+]] = OpLoad {{.*}} [[gep2]] -// CHECK: OpIAdd {{.*}} [[ld2]] [[ld1]] -// CHECK: OpFunctionEnd -// CHECK-NEXT: [[bar1]] = OpFunction -// CHECK-NOT: OpFunctionParameter -// CHECK: [[gep1:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var1]] [[int0]] [[int0]] -// CHECK: [[gep2:%[a-zA-Z0-9_]+]] = OpAccessChain {{.*}} [[var2]] [[int0]] [[int0]] -// CHECK: [[ld1:%[a-zA-Z0-9_]+]] = OpLoad {{.*}} [[gep1]] -// CHECK: [[ld2:%[a-zA-Z0-9_]+]] = OpLoad {{.*}} [[gep2]] -// CHECK: OpIAdd {{.*}} [[ld2]] [[ld1]] -// CHECK: OpFunctionEnd -// CHECK-NOT: OpFunction diff --git a/test/UBO/multiple_ubo_args.ll b/test/UBO/multiple_ubo_args.ll new file mode 100644 index 000000000..7938f5ebc --- /dev/null +++ b/test/UBO/multiple_ubo_args.ll @@ -0,0 +1,71 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; CHECK: define {{.*}} @k1 +; CHECK: call {{.*}} [[bar1:@bar[a-zA-Z0-9_.]+]]() +; CHECK: define {{.*}} @k2 +; CHECK: call {{.*}} [[bar2:@bar[a-zA-Z0-9_.]+]]() + +; CHECK: define {{.*}} [[bar2]]() +; CHECK-DAG: [[res1:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK-DAG: [[res2:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK-DAG: [[gep1:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res1]] +; CHECK-DAG: [[gep2:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res2]] +; CHECK-DAG: [[ld1:%[a-zA-Z0-9_.]+]] = load <4 x i32>, ptr addrspace(2) [[gep1]] +; CHECK-DAG: [[ld2:%[a-zA-Z0-9_.]+]] = load <4 x i32>, ptr addrspace(2) [[gep2]] + +; CHECK: define {{.*}} [[bar1]]() +; CHECK-DAG: [[res1:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK-DAG: [[res2:%[a-zA-Z0-9_.]+]] = call ptr addrspace(2) @_Z14clspv.resource +; CHECK-DAG: [[gep1:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res1]] +; CHECK-DAG: [[gep2:%[a-zA-Z0-9_.]+]] = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) [[res2]] +; CHECK-DAG: [[ld1:%[a-zA-Z0-9_.]+]] = load <4 x i32>, ptr addrspace(2) [[gep1]] +; CHECK-DAG: [[ld2:%[a-zA-Z0-9_.]+]] = load <4 x i32>, ptr addrspace(2) [[gep2]] + +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %in1, ptr addrspace(2) nocapture readonly %in2) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %in1, align 16 + %1 = load <4 x i32>, ptr addrspace(2) %in2, align 16 + %add = add <4 x i32> %1, %0 + ret <4 x i32> %add +} + +define dso_local spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in1, ptr addrspace(2) nocapture readonly align 16 %in2) !clspv.pod_args_impl !18 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = call ptr addrspace(2) @_Z14clspv.resource.2(i32 0, i32 2, i32 1, i32 2, i32 2, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %5 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %4, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %3, ptr addrspace(2) %5) #2 + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +define dso_local spir_kernel void @k2(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in1, ptr addrspace(2) nocapture readonly align 16 %in2) !clspv.pod_args_impl !18 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = call ptr addrspace(2) @_Z14clspv.resource.2(i32 0, i32 2, i32 1, i32 2, i32 2, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %5 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %4, i32 0, i32 0, i32 0 + %call = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %5, ptr addrspace(2) %3) #2 + store <4 x i32> %call, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +!18 = !{i32 2} + diff --git a/test/UBO/needs_inlined.cl b/test/UBO/needs_inlined.cl deleted file mode 100644 index 75c677c3c..000000000 --- a/test/UBO/needs_inlined.cl +++ /dev/null @@ -1,18 +0,0 @@ -// RUN: clspv %target %s -o %t.spv -constant-args-ubo -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm - -// This test is not validated because it uses a selection between Uniform -// pointers, which is disallowed by SPIR-V. -int4 bar(constant int4* data) { return data[0]; } - -kernel void k1(global int4* out, constant int4* in1, constant int4* in2, int a) { - constant int4* x = (a == 0) ? in1 : in2; - *out = bar(x); -} - -// CHECK: [[sel:%[a-zA-Z0-9_]+]] = OpSelect -// CHECK-NEXT: OpLoad {{.*}} [[sel]] -// CHECK-NOT: OpFunctionCall -// CHECK: OpFunctionEnd -// CHECK-NOT: OpFunction diff --git a/test/UBO/needs_inlined.ll b/test/UBO/needs_inlined.ll new file mode 100644 index 000000000..138fec63d --- /dev/null +++ b/test/UBO/needs_inlined.ll @@ -0,0 +1,55 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t.ll --passes=multi-version-ubo-functions,remove-unused-arguments +; RUN: FileCheck %s < %t.ll + +; This test would lead to invalid SPIR-V. + +; CHECK-NOT: define {{.*}} @bar +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 = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_func <4 x i32> @bar(ptr addrspace(2) nocapture readonly %data) { +entry: + %0 = load <4 x i32>, ptr addrspace(2) %data, align 16 + ret <4 x i32> %0 +} + +define spir_kernel void @k1(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in1, ptr addrspace(2) nocapture readonly align 16 %in2, { i32 } %podargs) !clspv.pod_args_impl !17 !kernel_arg_map !18 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x <4 x i32>] } zeroinitializer) + %1 = getelementptr { [0 x <4 x i32>] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0 + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %3 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %2, i32 0, i32 0, i32 0 + %4 = call ptr addrspace(2) @_Z14clspv.resource.2(i32 0, i32 2, i32 1, i32 2, i32 2, i32 0, { [4096 x <4 x i32>] } zeroinitializer) + %5 = getelementptr { [4096 x <4 x i32>] }, ptr addrspace(2) %4, i32 0, i32 0, i32 0 + %6 = call ptr addrspace(9) @_Z14clspv.resource.3(i32 -1, i32 3, i32 5, i32 3, i32 3, i32 0, { { i32 } } zeroinitializer) + %7 = getelementptr { { i32 } }, ptr addrspace(9) %6, i32 0, i32 0 + %8 = load { i32 }, ptr addrspace(9) %7, align 4 + %a = extractvalue { i32 } %8, 0 + %cmp.i = icmp eq i32 %a, 0 + %in1.in2 = select i1 %cmp.i, ptr addrspace(2) %3, ptr addrspace(2) %5 + %call.i = tail call spir_func <4 x i32> @bar(ptr addrspace(2) %in1.in2) + store <4 x i32> %call.i, ptr addrspace(1) %1, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +declare ptr addrspace(2) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { [4096 x <4 x i32>] }) + +declare ptr addrspace(9) @_Z14clspv.resource.3(i32, i32, i32, i32, i32, i32, { { i32 } }) + +attributes #0 = { nofree noinline norecurse nounwind memory(read) "no-builtins" "no-trapping-math"="true" "stack-protector-buffer-size"="0" "stackrealign" } +attributes #1 = { nofree norecurse nounwind memory(read, argmem: readwrite) "no-builtins" "no-trapping-math"="true" "stack-protector-buffer-size"="0" "stackrealign" "uniform-work-group-size"="true" } +attributes #2 = { nobuiltin nounwind "no-builtins" } + +!17 = !{i32 2} +!18 = !{!19, !20, !21, !22} +!19 = !{!"out", i32 0, i32 0, i32 0, i32 0, !"buffer"} +!20 = !{!"in1", i32 1, i32 1, i32 0, i32 0, !"buffer_ubo"} +!21 = !{!"in2", i32 2, i32 2, i32 0, i32 0, !"buffer_ubo"} +!22 = !{!"a", i32 3, i32 3, i32 0, i32 4, !"pod_pushconstant"} + diff --git a/test/UBO/nested_padding.cl b/test/UBO/nested_padding.cl deleted file mode 100644 index bdeb66274..000000000 --- a/test/UBO/nested_padding.cl +++ /dev/null @@ -1,56 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -typedef struct { - int x; - int y __attribute((aligned(16))); -} inner; - -typedef struct { - inner i[2]; -} outer; - -__kernel void foo(__global outer* data, __constant outer* c) { - unsigned gid = get_global_id(0); - data[gid].i[0].x = c[gid].i[0].x; - data[gid].i[0].y = c[gid].i[0].y; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[inner:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[inner]] 1 Offset 4 -// CHECK-DAG: OpMemberDecorate [[inner]] 2 Offset 16 -// CHECK-DAG: OpMemberDecorate [[inner]] 3 Offset 20 -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 64 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 64 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[inner]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] -// CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 -// CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[inner]] [[two]] -// CHECK-DAG: [[outer:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[outer]] -// CHECK-DAG: [[block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[block]] -// CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 -// CHECK-DAG: [[ubo_array]] = OpTypeArray [[outer]] [[int_1024]] -// CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] -// CHECK-DAG: [[c_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] -// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c]] = OpVariable [[c_ptr]] Uniform -// CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_ele_ptr]] [[c]] [[zero]] {{.*}} [[zero]] [[zero]] [[zero]] -// CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_gep]] -// CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] {{.*}} [[zero]] [[zero]] -// CHECK: OpStore [[data_gep]] [[c_load]] diff --git a/test/UBO/nested_padding.ll b/test/UBO/nested_padding.ll new file mode 100644 index 000000000..89de6a686 --- /dev/null +++ b/test/UBO/nested_padding.ll @@ -0,0 +1,74 @@ +; RUN: clspv-opt -constant-args-ubo -int8=0 %s -o %t -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[inner:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[inner]] 1 Offset 4 +; CHECK-DAG: OpMemberDecorate [[inner]] 2 Offset 16 +; CHECK-DAG: OpMemberDecorate [[inner]] 3 Offset 20 +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 64 +; CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[c]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 64 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[inner]] = OpTypeStruct [[int]] [[int]] [[int]] [[int]] +; CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 +; CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[inner]] [[two]] +; CHECK-DAG: [[outer:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[outer]] +; CHECK-DAG: [[block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] +; CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[block]] +; CHECK-DAG: [[int_1024:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1024 +; CHECK-DAG: [[ubo_array]] = OpTypeArray [[outer]] [[int_1024]] +; CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[c_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] +; CHECK-DAG: [[c_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] +; CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] +; CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[c]] = OpVariable [[c_ptr]] Uniform +; CHECK: [[c_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_ele_ptr]] [[c]] [[zero]] {{.*}} [[zero]] [[zero]] [[zero]] +; CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_gep]] +; CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] {{.*}} [[zero]] [[zero]] +; CHECK: OpStore [[data_gep]] [[c_load]] + +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" + +%struct.outer = type { [2 x %struct.inner] } +%struct.inner = type { i32, [12 x i8], i32, [12 x i8] } + +@__spirv_GlobalInvocationId = local_unnamed_addr addrspace(5) global <3 x i32> zeroinitializer +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c) !clspv.pod_args_impl !13 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.outer] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [1024 x %struct.outer] } zeroinitializer) + %2 = getelementptr <3 x i32>, ptr addrspace(5) @__spirv_GlobalInvocationId, i32 0, i32 0 + %3 = load i32, ptr addrspace(5) %2, align 16 + %4 = getelementptr { [1024 x %struct.outer] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 0 + %5 = load i32, ptr addrspace(2) %4, align 16 + %6 = getelementptr { [0 x %struct.outer] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 0 + store i32 %5, ptr addrspace(1) %6, align 16 + %7 = getelementptr { [1024 x %struct.outer] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 2 + %8 = load i32, ptr addrspace(2) %7, align 16 + %9 = getelementptr { [0 x %struct.outer] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %3, i32 0, i32 0, i32 2 + store i32 %8, ptr addrspace(1) %9, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.outer] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [1024 x %struct.outer] }) + +!13 = !{i32 2} + diff --git a/test/UBO/odd_size_padding.cl b/test/UBO/odd_size_padding.cl deleted file mode 100644 index 9c80fb4a7..000000000 --- a/test/UBO/odd_size_padding.cl +++ /dev/null @@ -1,33 +0,0 @@ -// RUN: clspv %target -int8 -constant-args-ubo -inline-entry-points %s -o %t.spv -// RUN: spirv-dis %t.spv -o %t.spvasm -// RUN: FileCheck %s < %t.spvasm -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// TODO(#1292) -// XFAIL: * - -typedef struct { - char a; - int b __attribute__((aligned(16))); -} S; - - kernel void foo(global S* out, constant S* in) { - out->b = in->b; -} - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 1 -// CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 16 -// CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 20 -// CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 -// CHECK: OpDecorate [[in]] NonWritable -// CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 32 -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 -// CHECK: [[s]] = OpTypeStruct [[char]] [[char]] [[int]] [[char]] -// CHECK-DAG: [[int_2048:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2048 -// CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_2048]] -// CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] -// CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform diff --git a/test/UBO/odd_size_padding.ll b/test/UBO/odd_size_padding.ll new file mode 100644 index 000000000..2dbce2d10 --- /dev/null +++ b/test/UBO/odd_size_padding.ll @@ -0,0 +1,45 @@ +; RUN: clspv-opt -constant-args-ubo %s -o %t -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm + +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 1 +; CHECK-DAG: OpMemberDecorate [[s]] 2 Offset 16 +; CHECK-DAG: OpMemberDecorate [[s]] 3 Offset 20 +; CHECK-DAG: OpDecorate [[in:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[in]] DescriptorSet 0 +; CHECK: OpDecorate [[in]] NonWritable +; CHECK: OpDecorate [[ubo_array:%[0-9a-zA-Z_]+]] ArrayStride 32 +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[char:%[0-9a-zA-Z_]+]] = OpTypeInt 8 0 +; CHECK: [[s]] = OpTypeStruct [[char]] [[char]] [[int]] [[char]] +; CHECK-DAG: [[int_2048:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2048 +; CHECK-DAG: [[ubo_array]] = OpTypeArray [[s]] [[int_2048]] +; CHECK-DAG: [[ubo_block:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[ubo_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_block]] +; CHECK-DAG: [[in]] = OpVariable [[ubo_ptr]] Uniform + +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" + +%struct.S = type { i8, [15 x i8], i32, [12 x i8] } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %out, ptr addrspace(2) nocapture readonly align 16 %in, { i32 } %podargs) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.S] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [2048 x %struct.S] } zeroinitializer) + %2 = getelementptr { [2048 x %struct.S] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0, i32 2 + %3 = load i32, ptr addrspace(2) %2, align 16 + %4 = getelementptr { [0 x %struct.S] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0, i32 2 + store i32 %3, ptr addrspace(1) %4, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.S] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [2048 x %struct.S] }) + +!14 = !{i32 2} diff --git a/test/UBO/test_cluster_pod_args.ll b/test/UBO/test_cluster_pod_args.ll new file mode 100644 index 000000000..a7d01283a --- /dev/null +++ b/test/UBO/test_cluster_pod_args.ll @@ -0,0 +1,51 @@ +; RUN: clspv-opt -constant-args-ubo -pod-ubo %s -o %t -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck %s < %t.map +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; TODO(#1303): invalid LLVM IR is produced in SPIRVProducer +; XFAIL: * + +; Just checking that the argument names are recorded correctly when clustering pod args. + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c_arg,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo +; MAP-NEXT: kernel,foo,arg,n,argOrdinal,2,descriptorSet,0,binding,2,offset,0,argKind,pod_ubo + +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" + +%struct.data_type = type { i32, [12 x i8] } + +@c_var = local_unnamed_addr addrspace(2) constant [2 x %struct.data_type] [%struct.data_type { i32 0, [12 x i8] undef }, %struct.data_type { i32 1, [12 x i8] undef }], align 16 +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c_arg, { i32 } %podargs) !clspv.pod_args_impl !8 !kernel_arg_map !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.data_type] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x %struct.data_type] } zeroinitializer) + %2 = call ptr addrspace(6) @_Z14clspv.resource.2(i32 0, i32 2, i32 4, i32 2, i32 2, i32 0, { { i32 } } zeroinitializer) + %3 = getelementptr { { i32 } }, ptr addrspace(6) %2, i32 0, i32 0 + %4 = load { i32 }, ptr addrspace(6) %3, align 4 + %n = extractvalue { i32 } %4, 0 + %5 = getelementptr { [4096 x %struct.data_type] }, ptr addrspace(2) %1, i32 0, i32 0, i32 %n, i32 0 + %6 = load i32, ptr addrspace(2) %5, align 16 + %7 = getelementptr [2 x %struct.data_type], ptr addrspace(2) @c_var, i32 0, i32 %n, i32 0 + %8 = load i32, ptr addrspace(2) %7, align 16 + %add.i = add nsw i32 %8, %6 + %9 = getelementptr { [0 x %struct.data_type] }, ptr addrspace(1) %0, i32 0, i32 0, i32 %n, i32 0 + store i32 %add.i, ptr addrspace(1) %9, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.data_type] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x %struct.data_type] }) + +declare ptr addrspace(6) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { { i32 } }) + +!8 = !{i32 1} +!14 = !{!15, !16, !17} +!15 = !{!"data", i32 0, i32 0, i32 0, i32 0, !"buffer"} +!16 = !{!"c_arg", i32 1, i32 1, i32 0, i32 0, !"buffer_ubo"} +!17 = !{!"n", i32 2, i32 2, i32 0, i32 4, !"pod_ubo"} + diff --git a/test/UBO/transform_local.cl b/test/UBO/transform_local.cl deleted file mode 100644 index ab0c0fc44..000000000 --- a/test/UBO/transform_local.cl +++ /dev/null @@ -1,71 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -arch=spir -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm --check-prefixes=CHECK,CHECK-32 -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -arch=spir64 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm --check-prefixes=CHECK,CHECK-64 -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv - -// TODO(#1292) -// XFAIL: * - -typedef struct { - int x __attribute__((aligned(16))); -} data_type; - -__kernel void foo(__global data_type* data, __constant data_type* c_arg, __local data_type* l_arg) { - data[2].x = c_arg[2].x + l_arg[2].x; -} - -// Most important thing here is the arrayElemSize check for the pointer-to-local arg. -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c_arg,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo -// MAP-NEXT: kernel,foo,arg,l_arg,argOrdinal,2,argKind,local,arrayElemSize,16,arrayNumElemSpecId,3 - -// CHECK-DAG: OpMemberDecorate [[data_type:%[0-9a-zA-Z_]+]] 1 Offset 4 -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c_arg:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c_arg]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[spec_id:%[0-9a-zA-Z_]+]] SpecId 3 -// CHECK-NOT: OpExtension -// CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK-DAG: [[data_type]] = OpTypeStruct [[int]] [[int]] -// CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[data_type]] -// CHECK-DAG: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] -// CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK-DAG: [[ubo_array:%[0-9a-zA-Z_]+]] = OpTypeArray [[data_type]] [[int_4096]] -// CHECK-DAG: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK-DAG: [[c_arg_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] -// CHECK-DAG: [[size1:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 -// CHECK-DAG: [[size2:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 -// CHECK-DAG: [[size3:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 -// CHECK-DAG: [[size:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 -// CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[data_type]] [[size]] -// CHECK-DAG: [[l_arg_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Workgroup [[array]] -// CHECK-DAG: [[c_arg_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] -// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 -// CHECK-DAG: [[l_arg_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Workgroup [[int]] -// CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] -// CHECK-64-DAG: [[long:%[0-9a-zA-Z_]+]] = OpTypeInt 64 0 -// CHECK-64-DAG: [[long_two:%[0-9a-zA-Z_]+]] = OpConstant [[long]] 2 -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c_arg]] = OpVariable [[c_arg_ptr]] Uniform -// CHECK: [[l_arg:%[0-9a-zA-Z_]+]] = OpVariable [[l_arg_ptr]] Workgroup -// CHECK: [[c_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_arg_ele_ptr]] [[c_arg]] [[zero]] [[two]] [[zero]] -// CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_arg_gep]] -// CHECK-64: [[l_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[l_arg_ele_ptr]] [[l_arg]] [[long_two]] [[zero]] -// CHECK-32: [[l_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[l_arg_ele_ptr]] [[l_arg]] [[two]] [[zero]] -// CHECK: [[l_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[l_arg_gep]] -// CHECK: [[add:%[0-9a-zA-Z_]+]] = OpIAdd [[int]] [[l_load]] [[c_load]] -// CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] [[two]] [[zero]] -// CHECK: OpStore [[data_gep]] [[add]] diff --git a/test/UBO/transform_local.ll b/test/UBO/transform_local.ll new file mode 100644 index 000000000..6b4f18dd7 --- /dev/null +++ b/test/UBO/transform_local.ll @@ -0,0 +1,90 @@ +; RUN: clspv-opt %s -o %t -constant-args-ubo -int8=0 -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; Most important thing here is the arrayElemSize check for the pointer-to-local arg. +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c_arg,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo +; MAP-NEXT: kernel,foo,arg,l_arg,argOrdinal,2,argKind,local,arrayElemSize,16,arrayNumElemSpecId,3 + +; CHECK-DAG: OpMemberDecorate [[data_type:%[0-9a-zA-Z_]+]] 1 Offset 4 +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c_arg:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[c_arg]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[spec_id:%[0-9a-zA-Z_]+]] SpecId 3 +; CHECK-NOT: OpExtension +; CHECK-DAG: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[data_type]] = OpTypeStruct [[int]] [[int]] +; CHECK-DAG: [[runtime]] = OpTypeRuntimeArray [[data_type]] +; CHECK-DAG: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] +; CHECK-DAG: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] +; CHECK-DAG: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK-DAG: [[ubo_array:%[0-9a-zA-Z_]+]] = OpTypeArray [[data_type]] [[int_4096]] +; CHECK-DAG: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK-DAG: [[c_arg_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] +; CHECK-DAG: [[size1:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 +; CHECK-DAG: [[size2:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 +; CHECK-DAG: [[size3:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 +; CHECK-DAG: [[size:%[0-9a-zA-Z_]+]] = OpSpecConstant [[int]] 1 +; CHECK-DAG: [[array:%[0-9a-zA-Z_]+]] = OpTypeArray [[data_type]] [[size]] +; CHECK-DAG: [[l_arg_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Workgroup [[array]] +; CHECK-DAG: [[c_arg_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] +; CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK-DAG: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 +; CHECK-DAG: [[l_arg_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Workgroup [[int]] +; CHECK-DAG: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] +; CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[c_arg]] = OpVariable [[c_arg_ptr]] Uniform +; CHECK: [[l_arg:%[0-9a-zA-Z_]+]] = OpVariable [[l_arg_ptr]] Workgroup +; CHECK: [[c_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_arg_ele_ptr]] [[c_arg]] [[zero]] [[two]] [[zero]] +; CHECK: [[c_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_arg_gep]] +; CHECK: [[l_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[l_arg_ele_ptr]] [[l_arg]] [[two]] [[zero]] +; CHECK: [[l_load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[l_arg_gep]] +; CHECK: [[add:%[0-9a-zA-Z_]+]] = OpIAdd [[int]] [[l_load]] [[c_load]] +; CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] [[two]] [[zero]] +; CHECK: OpStore [[data_gep]] [[add]] + +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" + +%struct.data_type = type { i32, [12 x i8] } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c_arg, ptr addrspace(3) nocapture readonly align 16 %l_arg) !clspv.pod_args_impl !17 { +entry: + %0 = call ptr addrspace(3) @_Z11clspv.local.3(i32 3, [0 x %struct.data_type] zeroinitializer) + %1 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.data_type] } zeroinitializer) + %2 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x %struct.data_type] } zeroinitializer) + %6 = getelementptr { [4096 x %struct.data_type] }, ptr addrspace(2) %2, i32 0, i32 0, i32 2, i32 0 + %7 = load i32, ptr addrspace(2) %6, align 16 + %8 = getelementptr [0 x %struct.data_type], ptr addrspace(3) %0, i32 0, i32 2, i32 0 + %9 = load i32, ptr addrspace(3) %8, align 16 + %add.i = add nsw i32 %9, %7 + %10 = getelementptr { [0 x %struct.data_type] }, ptr addrspace(1) %1, i32 0, i32 0, i32 2, i32 0 + store i32 %add.i, ptr addrspace(1) %10, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.data_type] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x %struct.data_type] }) + +declare ptr addrspace(9) @_Z14clspv.resource.2(i32, i32, i32, i32, i32, i32, { { i32 } }) + +declare ptr addrspace(3) @_Z11clspv.local.3(i32, [0 x %struct.data_type]) + +!clspv.next_spec_constant_id = !{!9} +!clspv.spec_constant_list = !{!10} +!_Z20clspv.local_spec_ids = !{!11} + +!9 = distinct !{i32 4} +!10 = !{i32 3, i32 3} +!11 = !{ptr @foo, i32 2, i32 3} +!17 = !{i32 2} + diff --git a/test/UBO/transform_padding.cl b/test/UBO/transform_padding.cl deleted file mode 100644 index 28de9e820..000000000 --- a/test/UBO/transform_padding.cl +++ /dev/null @@ -1,49 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -int8=0 -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv -// TODO(#1292) -// XFAIL: * - -// The data_type struct translates as { i32, [12 x i8] } which is transformed -// to { i32, i32 } -typedef struct { - int x __attribute((aligned(16))); -} data_type; - -__kernel void foo(__global data_type *data, __constant data_type *c_arg) { - data[2].x = c_arg[2].x; -} - -// MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c_arg,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK-DAG: OpMemberDecorate [[data_type:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[data_type]] 1 Offset 4 -// CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 -// CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c_arg:%[0-9a-zA-Z_]+]] Binding 1 -// CHECK-DAG: OpDecorate [[c_arg]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[c_arg]] NonWritable -// CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK: [[data_type]] = OpTypeStruct [[int]] [[int]] -// CHECK: [[runtime]] = OpTypeRuntimeArray [[data_type]] -// CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] -// CHECK: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] -// CHECK: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK: [[ubo_array:%[0-9a-zA-Z_]+]] = OpTypeArray [[data_type]] [[int_4096]] -// CHECK: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] -// CHECK: [[c_arg_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] -// CHECK: [[c_arg_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] -// CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 -// CHECK: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] -// CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer -// CHECK: [[c_arg]] = OpVariable [[c_arg_ptr]] Uniform -// CHECK: [[c_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_arg_ele_ptr]] [[c_arg]] [[zero]] [[two]] [[zero]] -// CHECK: [[load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_arg_gep]] -// CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] [[two]] [[zero]] -// CHECK: OpStore [[data_gep]] [[load]] diff --git a/test/UBO/transform_padding.ll b/test/UBO/transform_padding.ll new file mode 100644 index 000000000..73d25facf --- /dev/null +++ b/test/UBO/transform_padding.ll @@ -0,0 +1,62 @@ +; RUN: clspv-opt %s -o %t -constant-args-ubo -int8=0 -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,data,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c_arg,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpDecorate [[runtime:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK-DAG: OpMemberDecorate [[data_type:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[data_type]] 1 Offset 4 +; CHECK-DAG: OpDecorate [[data:%[0-9a-zA-Z_]+]] Binding 0 +; CHECK-DAG: OpDecorate [[data]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c_arg:%[0-9a-zA-Z_]+]] Binding 1 +; CHECK-DAG: OpDecorate [[c_arg]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[c_arg]] NonWritable +; CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK: [[data_type]] = OpTypeStruct [[int]] [[int]] +; CHECK: [[runtime]] = OpTypeRuntimeArray [[data_type]] +; CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[runtime]] +; CHECK: [[data_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[struct]] +; CHECK: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK: [[ubo_array:%[0-9a-zA-Z_]+]] = OpTypeArray [[data_type]] [[int_4096]] +; CHECK: [[ubo_struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[ubo_array]] +; CHECK: [[c_arg_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[ubo_struct]] +; CHECK: [[c_arg_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int]] +; CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK: [[two:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 2 +; CHECK: [[data_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer StorageBuffer [[int]] +; CHECK: [[data]] = OpVariable [[data_ptr]] StorageBuffer +; CHECK: [[c_arg]] = OpVariable [[c_arg_ptr]] Uniform +; CHECK: [[c_arg_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[c_arg_ele_ptr]] [[c_arg]] [[zero]] [[two]] [[zero]] +; CHECK: [[load:%[0-9a-zA-Z_]+]] = OpLoad [[int]] [[c_arg_gep]] +; CHECK: [[data_gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[data_ele_ptr]] [[data]] [[zero]] [[two]] [[zero]] +; CHECK: OpStore [[data_gep]] [[load]] + +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" + +%struct.data_type = type { i32, [12 x i8] } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 16 %data, ptr addrspace(2) nocapture readonly align 16 %c_arg, { i32 } %podargs) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.data_type] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x %struct.data_type] } zeroinitializer) + %5 = getelementptr { [4096 x %struct.data_type] }, ptr addrspace(2) %1, i32 0, i32 0, i32 2, i32 0 + %6 = load i32, ptr addrspace(2) %5, align 16 + %7 = getelementptr { [0 x %struct.data_type] }, ptr addrspace(1) %0, i32 0, i32 0, i32 2, i32 0 + store i32 %6, ptr addrspace(1) %7, align 16 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.data_type] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x %struct.data_type] }) + +!14 = !{i32 2} + diff --git a/test/UBO/vec2_no_pad.cl b/test/UBO/vec2_no_pad.cl deleted file mode 100644 index 8a998374d..000000000 --- a/test/UBO/vec2_no_pad.cl +++ /dev/null @@ -1,42 +0,0 @@ -// RUN: clspv %target -constant-args-ubo -inline-entry-points %s -o %t.spv -// RUN: spirv-dis -o %t2.spvasm %t.spv -// RUN: FileCheck %s < %t2.spvasm -// RUN: clspv-reflection %t.spv -o %t2.map -// RUN: FileCheck -check-prefix=MAP %s < %t2.map -// RUN: spirv-val --target-env vulkan1.0 %t.spv -// TODO(#1292) -// XFAIL: * - -// Natural alignment don't lead to LLVM inserting packing so this is ok. -typedef struct { - int x; - int2 y; -} data_type; - -__kernel void foo(__global data_type* d, __constant data_type* c) { - d->y = c->y; -} - -// MAP: kernel,foo,arg,d,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer -// MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo - -// CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 8 -// CHECK-DAG: OpMemberDecorate [[struct:%[0-9a-zA-Z_]+]] 0 Offset 0 -// CHECK-DAG: OpDecorate [[var:%[0-9a-zA-Z_]+]] NonWritable -// CHECK-DAG: OpDecorate [[var]] DescriptorSet 0 -// CHECK-DAG: OpDecorate [[var]] Binding 1 -// CHECK: OpDecorate [[array:%[0-9a-zA-Z_]+]] ArrayStride 16 -// CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 -// CHECK: [[int2:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 2 -// CHECK: [[s]] = OpTypeStruct [[int]] [[int2]] -// CHECK: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 -// CHECK: [[array]] = OpTypeArray [[s]] [[int_4096]] -// CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] -// CHECK: [[ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct]] -// CHECK: [[ptr_int2:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int2]] -// CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 -// CHECK: [[one:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1 -// CHECK: [[var]] = OpVariable [[ptr]] Uniform -// CHECK: [[gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_int2]] [[var]] [[zero]] [[zero]] [[one]] -// CHECK: OpLoad [[int2]] [[gep]] diff --git a/test/UBO/vec2_no_pad.ll b/test/UBO/vec2_no_pad.ll new file mode 100644 index 000000000..c43f0d5df --- /dev/null +++ b/test/UBO/vec2_no_pad.ll @@ -0,0 +1,55 @@ +; RUN: clspv-opt %s -o %t -constant-args-ubo -producer-out-file %t.spv --passes=ubo-type-transform,spirv-producer +; RUN: spirv-dis %t.spv -o %t.spvasm +; RUN: spirv-val --target-env vulkan1.0 %t.spv +; RUN: FileCheck %s < %t.spvasm +; RUN: clspv-reflection %t.spv -o %t.map +; RUN: FileCheck --check-prefix=MAP %s < %t.map + +; MAP: kernel,foo,arg,d,argOrdinal,0,descriptorSet,0,binding,0,offset,0,argKind,buffer +; MAP-NEXT: kernel,foo,arg,c,argOrdinal,1,descriptorSet,0,binding,1,offset,0,argKind,buffer_ubo + +; CHECK-DAG: OpMemberDecorate [[s:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpMemberDecorate [[s]] 1 Offset 8 +; CHECK-DAG: OpMemberDecorate [[struct:%[0-9a-zA-Z_]+]] 0 Offset 0 +; CHECK-DAG: OpDecorate [[var:%[0-9a-zA-Z_]+]] NonWritable +; CHECK-DAG: OpDecorate [[var]] DescriptorSet 0 +; CHECK-DAG: OpDecorate [[var]] Binding 1 +; CHECK: OpDecorate [[array:%[0-9a-zA-Z_]+]] ArrayStride 16 +; CHECK: [[int:%[0-9a-zA-Z_]+]] = OpTypeInt 32 0 +; CHECK: [[int2:%[0-9a-zA-Z_]+]] = OpTypeVector [[int]] 2 +; CHECK: [[s]] = OpTypeStruct [[int]] [[int2]] +; CHECK: [[int_4096:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 4096 +; CHECK: [[array]] = OpTypeArray [[s]] [[int_4096]] +; CHECK: [[struct:%[0-9a-zA-Z_]+]] = OpTypeStruct [[array]] +; CHECK: [[ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[struct]] +; CHECK: [[ptr_int2:%[0-9a-zA-Z_]+]] = OpTypePointer Uniform [[int2]] +; CHECK: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0 +; CHECK: [[one:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1 +; CHECK: [[var]] = OpVariable [[ptr]] Uniform +; CHECK: [[gep:%[0-9a-zA-Z_]+]] = OpAccessChain [[ptr_int2]] [[var]] [[zero]] [[zero]] [[one]] +; CHECK: OpLoad [[int2]] [[gep]] + +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" + +%struct.data_type = type { i32, <2 x i32> } + +@__spirv_WorkgroupSize = local_unnamed_addr addrspace(8) global <3 x i32> zeroinitializer + +define spir_kernel void @foo(ptr addrspace(1) nocapture writeonly align 8 %d, ptr addrspace(2) nocapture readonly align 8 %c) !clspv.pod_args_impl !14 { +entry: + %0 = call ptr addrspace(1) @_Z14clspv.resource.0(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, { [0 x %struct.data_type] } zeroinitializer) + %1 = call ptr addrspace(2) @_Z14clspv.resource.1(i32 0, i32 1, i32 1, i32 1, i32 1, i32 0, { [4096 x %struct.data_type] } zeroinitializer) + %5 = getelementptr { [4096 x %struct.data_type] }, ptr addrspace(2) %1, i32 0, i32 0, i32 0, i32 1 + %6 = load <2 x i32>, ptr addrspace(2) %5, align 8 + %7 = getelementptr { [0 x %struct.data_type] }, ptr addrspace(1) %0, i32 0, i32 0, i32 0, i32 1 + store <2 x i32> %6, ptr addrspace(1) %7, align 8 + ret void +} + +declare ptr addrspace(1) @_Z14clspv.resource.0(i32, i32, i32, i32, i32, i32, { [0 x %struct.data_type] }) + +declare ptr addrspace(2) @_Z14clspv.resource.1(i32, i32, i32, i32, i32, i32, { [4096 x %struct.data_type] }) + +!14 = !{i32 2} +