From dade7201cdaa1e493db48f6d57f48e6a51ee1dc4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?K=C3=A9vin=20Petit?= Date: Sun, 30 Oct 2022 18:54:16 +0000 Subject: [PATCH 1/2] WIP texel buffers --- include/clspv/ArgKind.h | 2 ++ lib/AllocateDescriptorsPass.cpp | 9 +++++++-- lib/ArgKind.cpp | 8 ++++++++ lib/SPIRVProducerPass.cpp | 20 +++++++++++++++++-- lib/Types.cpp | 10 ++++++++++ lib/Types.h | 2 ++ .../readwrite_texel_buffer_argument.cl | 17 ++++++++++++++++ .../storage_texel_buffer_argument.cl | 19 ++++++++++++++++++ .../uniform_texel_buffer_argument.cl | 19 ++++++++++++++++++ 9 files changed, 102 insertions(+), 4 deletions(-) create mode 100644 test/Reflection/readwrite_texel_buffer_argument.cl create mode 100644 test/Reflection/storage_texel_buffer_argument.cl create mode 100644 test/Reflection/uniform_texel_buffer_argument.cl diff --git a/include/clspv/ArgKind.h b/include/clspv/ArgKind.h index b2c795847..fa0eae01f 100644 --- a/include/clspv/ArgKind.h +++ b/include/clspv/ArgKind.h @@ -31,6 +31,8 @@ enum class ArgKind : int { Sampler, PointerUBO, PointerPushConstant, + StorageTexelBuffer, + UniformTexelBuffer, }; // Converts an ArgKind to its string name. diff --git a/lib/AllocateDescriptorsPass.cpp b/lib/AllocateDescriptorsPass.cpp index fc6138d6b..c9925940b 100644 --- a/lib/AllocateDescriptorsPass.cpp +++ b/lib/AllocateDescriptorsPass.cpp @@ -281,9 +281,10 @@ bool clspv::AllocateDescriptorsPass::AllocateKernelArgDescriptors(Module &M) { int coherent = 0; if (uses_barriers && (arg_kind == clspv::ArgKind::Buffer || - arg_kind == clspv::ArgKind::StorageImage)) { + arg_kind == clspv::ArgKind::StorageImage || + arg_kind == clspv::ArgKind::StorageTexelBuffer)) { // Coherency is only required if the argument is an SSBO or storage - // image that is both read and written to. + // image or texel buffer that is both read and written to. bool reads = false; bool writes = false; std::tie(reads, writes) = HasReadsAndWrites(&Arg); @@ -584,6 +585,8 @@ bool clspv::AllocateDescriptorsPass::AllocateKernelArgDescriptors(Module &M) { case clspv::ArgKind::Sampler: case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: // We won't be translating the value here. Keep the type the same. // since calls using these values need to keep the same type. resource_type = inferred_ty; @@ -665,6 +668,8 @@ bool clspv::AllocateDescriptorsPass::AllocateKernelArgDescriptors(Module &M) { } break; case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: case clspv::ArgKind::Sampler: { // The call returns a pointer to an opaque type. Eventually the // SPIR-V will need to load the variable, so the natural thing would diff --git a/lib/ArgKind.cpp b/lib/ArgKind.cpp index ee72051f7..3acb9ca5d 100644 --- a/lib/ArgKind.cpp +++ b/lib/ArgKind.cpp @@ -165,6 +165,10 @@ const char *GetArgKindName(ArgKind kind) { return "pointer_pushconstant"; case ArgKind::PointerUBO: return "pointer_ubo"; + case ArgKind::StorageTexelBuffer: + return "storage_texel_buffer"; + case ArgKind::UniformTexelBuffer: + return "uniform_texel_buffer"; } errs() << "Unhandled case in clspv::GetArgKindForType: " << int(kind) << "\n"; llvm_unreachable("Unhandled case in clspv::GetArgKindForType"); @@ -194,6 +198,10 @@ ArgKind GetArgKindFromName(const std::string &name) { return ArgKind::PointerPushConstant; } else if (name == "pointer_ubo") { return ArgKind::PointerUBO; + } else if (name == "storage_texel_buffer") { + return ArgKind::StorageTexelBuffer; + } else if (name == "uniform_texel_buffer") { + return ArgKind::UniformTexelBuffer; } llvm_unreachable("Unhandled case in clspv::GetArgKindFromName"); return ArgKind::Buffer; diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index d96c241a1..a99c4207d 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -1266,6 +1266,8 @@ void SPIRVProducerPassImpl::FindTypesForResourceVars() { case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: case clspv::ArgKind::Sampler: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: break; default: break; @@ -1498,6 +1500,8 @@ spv::StorageClass SPIRVProducerPassImpl::GetStorageClassForArgKind( case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: case clspv::ArgKind::Sampler: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: return spv::StorageClassUniformConstant; default: llvm_unreachable("Unsupported storage class for argument kind"); @@ -1930,7 +1934,7 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { RID = addSPIRVInst(spv::OpTypeImage, Ops); // Only need a sampled version of the type if it is used with a sampler. - if (Sampled == 1 && ImageDimensionality(ext_ty) != spv::DimBuffer) { + if (Sampled == 1 && ImageDimensionality(ext_ty) != spv::DimBuffer) { // TODO why is DimBuffer excluded? Ops.clear(); Ops << RID; getImageTypeMap()[Canonical] = @@ -2516,6 +2520,8 @@ void SPIRVProducerPassImpl::GenerateResourceVars() { case clspv::ArgKind::Sampler: case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: type = info->data_type->getPointerTo(AddressSpace::UniformConstant); break; default: @@ -2547,6 +2553,8 @@ void SPIRVProducerPassImpl::GenerateResourceVars() { case clspv::ArgKind::Sampler: case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: // The call maps to a load we generate later. ResourceVarDeferredLoadCalls[call] = info->var_id; break; @@ -3745,7 +3753,7 @@ SPIRVProducerPassImpl::GenerateImageInstruction(CallInst *Call, Ops << Call->getType() << RID; RID = addSPIRVInst(spv::OpBitcast, Ops); } - } else if (IsStorageImageType(image_ty)) { + } else if (IsStorageImageType(image_ty) || IsStorageTexelBufferImageType(image_ty)) { // read_image on a storage image is mapped to OpImageRead. Value *Image = Call->getArgOperand(0); Value *Coordinate = Call->getArgOperand(1); @@ -7164,6 +7172,12 @@ void SPIRVProducerPassImpl::AddArgumentReflection( case clspv::ArgKind::Sampler: ext_inst = reflection::ExtInstArgumentSampler; break; + case clspv::ArgKind::StorageTexelBuffer: + ext_inst = reflection::ExtInstArgumentStorageTexelBuffer; + break; + case clspv::ArgKind::UniformTexelBuffer: + ext_inst = reflection::ExtInstArgumentUniformTexelBuffer; + break; default: llvm_unreachable("Unhandled argument reflection"); break; @@ -7180,6 +7194,8 @@ void SPIRVProducerPassImpl::AddArgumentReflection( case clspv::ArgKind::SampledImage: case clspv::ArgKind::StorageImage: case clspv::ArgKind::Sampler: + case clspv::ArgKind::StorageTexelBuffer: + case clspv::ArgKind::UniformTexelBuffer: Ops << getSPIRVInt32Constant(descriptor_set) << getSPIRVInt32Constant(binding); break; diff --git a/lib/Types.cpp b/lib/Types.cpp index a5f89ce56..b9ac4c75b 100644 --- a/lib/Types.cpp +++ b/lib/Types.cpp @@ -569,6 +569,16 @@ bool clspv::IsArrayImageType(llvm::Type *type) { return false; } +bool clspv::IsStorageTexelBufferImageType(llvm::StructType *type) { + if (ImageDimensionality(type) != spv::DimBuffer) + return false; + if (type->getName().contains("_wo") || + type->getName().contains("_rw")) { + return true; + } + return false; +} + bool clspv::IsFloatImageType(llvm::Type *type) { return IsImageType(type) && !IsIntImageType(type) && !IsUintImageType(type); } diff --git a/lib/Types.h b/lib/Types.h index 4a853003a..cdee62d13 100644 --- a/lib/Types.h +++ b/lib/Types.h @@ -64,6 +64,8 @@ bool IsSampledImageType(llvm::Type *type); // for read_write and write_only images. bool IsStorageImageType(llvm::Type *type); +bool IsStorageTexelBufferImageType(llvm::StructType *type); + // Returns true if the given type is a float image type. // Before image specialization, all images are considered float images. bool IsFloatImageType(llvm::Type *type); diff --git a/test/Reflection/readwrite_texel_buffer_argument.cl b/test/Reflection/readwrite_texel_buffer_argument.cl new file mode 100644 index 000000000..78bfb40fc --- /dev/null +++ b/test/Reflection/readwrite_texel_buffer_argument.cl @@ -0,0 +1,17 @@ +// RUN: clspv %s -o %t.spv -cl-std=CL2.0 -inline-entry-points +// RUN: spirv-dis %t.spv -o %t.spvasm +// RUN: FileCheck %s < %t.spvasm +// RUN: spirv-val %t.spv --target-env vulkan1.0 + +kernel void foo(read_write image1d_buffer_t im) { } + +// CHECK: [[import:%[a-zA-Z0-9_]+]] = OpExtInstImport "NonSemantic.ClspvReflection.2" +// CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] "foo" +// CHECK-DAG: [[foo_name:%[a-zA-Z0-9_]+]] = OpString "foo" +// CHECK-DAG: [[im_name:%[a-zA-Z0-9_]+]] = OpString "im" +// CHECK-DAG: [[void:%[a-zA-Z0-9_]+]] = OpTypeVoid +// CHECK-DAG: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 +// CHECK-DAG: [[int_0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 +// CHECK: [[kernel:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] Kernel [[foo]] [[foo_name]] +// CHECK: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentInfo [[im_name]] +// CHECK: [[arg:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentStorageTexelBuffer [[kernel]] [[int_0]] [[int_0]] [[int_0]] [[info]] diff --git a/test/Reflection/storage_texel_buffer_argument.cl b/test/Reflection/storage_texel_buffer_argument.cl new file mode 100644 index 000000000..8e576eb0f --- /dev/null +++ b/test/Reflection/storage_texel_buffer_argument.cl @@ -0,0 +1,19 @@ +// RUN: clspv %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 + +kernel void foo(write_only image1d_buffer_t data) { +} + +// CHECK: [[import:%[a-zA-Z0-9_]+]] = OpExtInstImport "NonSemantic.ClspvReflection.2" +// CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] "foo" +// CHECK-DAG: [[foo_name:%[a-zA-Z0-9_]+]] = OpString "foo" +// CHECK-DAG: [[data_name:%[a-zA-Z0-9_]+]] = OpString "data" +// CHECK-DAG: [[void:%[a-zA-Z0-9_]+]] = OpTypeVoid +// CHECK-DAG: [[uint:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 +// CHECK: [[uint_0:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 0 +// CHECK-DAG: [[decl:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] Kernel [[foo]] [[foo_name]] +// CHECK-DAG: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentInfo [[data_name]] +// CHECK: OpExtInst [[void]] [[import]] ArgumentStorageTexelBuffer [[decl]] [[uint_0]] [[uint_0]] [[uint_0]] [[info]] + diff --git a/test/Reflection/uniform_texel_buffer_argument.cl b/test/Reflection/uniform_texel_buffer_argument.cl new file mode 100644 index 000000000..fbb3a6f73 --- /dev/null +++ b/test/Reflection/uniform_texel_buffer_argument.cl @@ -0,0 +1,19 @@ +// RUN: clspv %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 + +kernel void foo(read_only image1d_buffer_t data) { +} + +// CHECK: [[import:%[a-zA-Z0-9_]+]] = OpExtInstImport "NonSemantic.ClspvReflection.2" +// CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] "foo" +// CHECK-DAG: [[foo_name:%[a-zA-Z0-9_]+]] = OpString "foo" +// CHECK-DAG: [[data_name:%[a-zA-Z0-9_]+]] = OpString "data" +// CHECK-DAG: [[void:%[a-zA-Z0-9_]+]] = OpTypeVoid +// CHECK-DAG: [[uint:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 +// CHECK: [[uint_0:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 0 +// CHECK-DAG: [[decl:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] Kernel [[foo]] [[foo_name]] +// CHECK-DAG: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentInfo [[data_name]] +// CHECK: OpExtInst [[void]] [[import]] ArgumentUniformTexelBuffer [[decl]] [[uint_0]] [[uint_0]] [[uint_0]] [[info]] + From 4e5a48d5ed287b245c10cfe8bfe4366478e23572 Mon Sep 17 00:00:00 2001 From: Romaric Jodin Date: Mon, 31 Jul 2023 14:27:03 +0200 Subject: [PATCH 2/2] fix after rebase --- lib/ArgKind.cpp | 8 ++++++-- lib/SPIRVProducerPass.cpp | 5 +++-- lib/Types.cpp | 10 ---------- lib/Types.h | 2 -- test/AllocateDescriptors/images.ll | 2 +- test/Reflection/readwrite_texel_buffer_argument.cl | 9 ++++----- test/Reflection/storage_texel_buffer_argument.cl | 10 ++++------ test/Reflection/uniform_texel_buffer_argument.cl | 9 ++++----- 8 files changed, 22 insertions(+), 33 deletions(-) diff --git a/lib/ArgKind.cpp b/lib/ArgKind.cpp index 3acb9ca5d..881e3ab16 100644 --- a/lib/ArgKind.cpp +++ b/lib/ArgKind.cpp @@ -57,9 +57,13 @@ clspv::ArgKind GetArgKindForType(Type *type) { return clspv::ArgKind::Sampler; if (clspv::IsImageType(ext_ty)) { if (clspv::IsSampledImageType(ext_ty)) { - return clspv::ArgKind::SampledImage; + return clspv::ImageDimensionality(ext_ty) == spv::DimBuffer + ? clspv::ArgKind::UniformTexelBuffer + : clspv::ArgKind::SampledImage; } else { - return clspv::ArgKind::StorageImage; + return clspv::ImageDimensionality(ext_ty) == spv::DimBuffer + ? clspv::ArgKind::StorageTexelBuffer + : clspv::ArgKind::StorageImage; } } errs() << "Unhandled target ext type: " << *type << "\n"; diff --git a/lib/SPIRVProducerPass.cpp b/lib/SPIRVProducerPass.cpp index a99c4207d..b1390a394 100644 --- a/lib/SPIRVProducerPass.cpp +++ b/lib/SPIRVProducerPass.cpp @@ -1934,7 +1934,8 @@ SPIRVID SPIRVProducerPassImpl::getSPIRVType(Type *Ty, bool needs_layout) { RID = addSPIRVInst(spv::OpTypeImage, Ops); // Only need a sampled version of the type if it is used with a sampler. - if (Sampled == 1 && ImageDimensionality(ext_ty) != spv::DimBuffer) { // TODO why is DimBuffer excluded? + // In SPIR-V 1.6 or later, sampled image dimension must not be Buffer + if (Sampled == 1 && ImageDimensionality(ext_ty) != spv::DimBuffer) { Ops.clear(); Ops << RID; getImageTypeMap()[Canonical] = @@ -3753,7 +3754,7 @@ SPIRVProducerPassImpl::GenerateImageInstruction(CallInst *Call, Ops << Call->getType() << RID; RID = addSPIRVInst(spv::OpBitcast, Ops); } - } else if (IsStorageImageType(image_ty) || IsStorageTexelBufferImageType(image_ty)) { + } else if (IsStorageImageType(image_ty)) { // read_image on a storage image is mapped to OpImageRead. Value *Image = Call->getArgOperand(0); Value *Coordinate = Call->getArgOperand(1); diff --git a/lib/Types.cpp b/lib/Types.cpp index b9ac4c75b..a5f89ce56 100644 --- a/lib/Types.cpp +++ b/lib/Types.cpp @@ -569,16 +569,6 @@ bool clspv::IsArrayImageType(llvm::Type *type) { return false; } -bool clspv::IsStorageTexelBufferImageType(llvm::StructType *type) { - if (ImageDimensionality(type) != spv::DimBuffer) - return false; - if (type->getName().contains("_wo") || - type->getName().contains("_rw")) { - return true; - } - return false; -} - bool clspv::IsFloatImageType(llvm::Type *type) { return IsImageType(type) && !IsIntImageType(type) && !IsUintImageType(type); } diff --git a/lib/Types.h b/lib/Types.h index cdee62d13..4a853003a 100644 --- a/lib/Types.h +++ b/lib/Types.h @@ -64,8 +64,6 @@ bool IsSampledImageType(llvm::Type *type); // for read_write and write_only images. bool IsStorageImageType(llvm::Type *type); -bool IsStorageTexelBufferImageType(llvm::StructType *type); - // Returns true if the given type is a float image type. // Before image specialization, all images are considered float images. bool IsFloatImageType(llvm::Type *type); diff --git a/test/AllocateDescriptors/images.ll b/test/AllocateDescriptors/images.ll index a0a5cdf7f..fdda420e6 100644 --- a/test/AllocateDescriptors/images.ll +++ b/test/AllocateDescriptors/images.ll @@ -49,7 +49,7 @@ entry: } ; CHECK-LABEL: @test4 -; CHECK: call [[image1dbufferf]] @_Z14clspv.resource.6(i32 1, i32 0, i32 7, i32 0, i32 6, i32 0, [[image1dbufferf]] +; CHECK: call [[image1dbufferf]] @_Z14clspv.resource.6(i32 1, i32 0, i32 11, i32 0, i32 6, i32 0, [[image1dbufferf]] define spir_kernel void @test4(target("spirv.Image", float, 5, 0, 0, 0, 2, 0, 1, 0) %im, ptr addrspace(1) nocapture readonly align 16 %data) !clspv.pod_args_impl !10 { entry: %0 = load <4 x float>, ptr addrspace(1) %data, align 16 diff --git a/test/Reflection/readwrite_texel_buffer_argument.cl b/test/Reflection/readwrite_texel_buffer_argument.cl index 78bfb40fc..3efb908b2 100644 --- a/test/Reflection/readwrite_texel_buffer_argument.cl +++ b/test/Reflection/readwrite_texel_buffer_argument.cl @@ -1,17 +1,16 @@ -// RUN: clspv %s -o %t.spv -cl-std=CL2.0 -inline-entry-points +// RUN: clspv %s -o %t.spv -cl-std=CL2.0 -inline-entry-points --cl-kernel-arg-info // RUN: spirv-dis %t.spv -o %t.spvasm // RUN: FileCheck %s < %t.spvasm // RUN: spirv-val %t.spv --target-env vulkan1.0 kernel void foo(read_write image1d_buffer_t im) { } -// CHECK: [[import:%[a-zA-Z0-9_]+]] = OpExtInstImport "NonSemantic.ClspvReflection.2" // CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] "foo" // CHECK-DAG: [[foo_name:%[a-zA-Z0-9_]+]] = OpString "foo" // CHECK-DAG: [[im_name:%[a-zA-Z0-9_]+]] = OpString "im" // CHECK-DAG: [[void:%[a-zA-Z0-9_]+]] = OpTypeVoid // CHECK-DAG: [[int:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 // CHECK-DAG: [[int_0:%[a-zA-Z0-9_]+]] = OpConstant [[int]] 0 -// CHECK: [[kernel:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] Kernel [[foo]] [[foo_name]] -// CHECK: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentInfo [[im_name]] -// CHECK: [[arg:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentStorageTexelBuffer [[kernel]] [[int_0]] [[int_0]] [[int_0]] [[info]] +// CHECK: [[kernel:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} Kernel [[foo]] [[foo_name]] +// CHECK: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} ArgumentInfo [[im_name]] +// CHECK: [[arg:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} ArgumentStorageTexelBuffer [[kernel]] [[int_0]] [[int_0]] [[int_0]] [[info]] diff --git a/test/Reflection/storage_texel_buffer_argument.cl b/test/Reflection/storage_texel_buffer_argument.cl index 8e576eb0f..134b7bc02 100644 --- a/test/Reflection/storage_texel_buffer_argument.cl +++ b/test/Reflection/storage_texel_buffer_argument.cl @@ -1,4 +1,4 @@ -// RUN: clspv %s -o %t.spv +// RUN: clspv %s -o %t.spv --cl-kernel-arg-info // RUN: spirv-dis %t.spv -o %t.spvasm // RUN: FileCheck %s < %t.spvasm // RUN: spirv-val --target-env vulkan1.0 %t.spv @@ -6,14 +6,12 @@ kernel void foo(write_only image1d_buffer_t data) { } -// CHECK: [[import:%[a-zA-Z0-9_]+]] = OpExtInstImport "NonSemantic.ClspvReflection.2" // CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] "foo" // CHECK-DAG: [[foo_name:%[a-zA-Z0-9_]+]] = OpString "foo" // CHECK-DAG: [[data_name:%[a-zA-Z0-9_]+]] = OpString "data" // CHECK-DAG: [[void:%[a-zA-Z0-9_]+]] = OpTypeVoid // CHECK-DAG: [[uint:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 // CHECK: [[uint_0:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 0 -// CHECK-DAG: [[decl:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] Kernel [[foo]] [[foo_name]] -// CHECK-DAG: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentInfo [[data_name]] -// CHECK: OpExtInst [[void]] [[import]] ArgumentStorageTexelBuffer [[decl]] [[uint_0]] [[uint_0]] [[uint_0]] [[info]] - +// CHECK-DAG: [[decl:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} Kernel [[foo]] [[foo_name]] +// CHECK-DAG: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} ArgumentInfo [[data_name]] +// CHECK: OpExtInst [[void]] {{.*}} ArgumentStorageTexelBuffer [[decl]] [[uint_0]] [[uint_0]] [[uint_0]] [[info]] diff --git a/test/Reflection/uniform_texel_buffer_argument.cl b/test/Reflection/uniform_texel_buffer_argument.cl index fbb3a6f73..46de30453 100644 --- a/test/Reflection/uniform_texel_buffer_argument.cl +++ b/test/Reflection/uniform_texel_buffer_argument.cl @@ -1,4 +1,4 @@ -// RUN: clspv %s -o %t.spv +// RUN: clspv %s -o %t.spv --cl-kernel-arg-info // RUN: spirv-dis %t.spv -o %t.spvasm // RUN: FileCheck %s < %t.spvasm // RUN: spirv-val --target-env vulkan1.0 %t.spv @@ -6,14 +6,13 @@ kernel void foo(read_only image1d_buffer_t data) { } -// CHECK: [[import:%[a-zA-Z0-9_]+]] = OpExtInstImport "NonSemantic.ClspvReflection.2" // CHECK: OpEntryPoint GLCompute [[foo:%[a-zA-Z0-9_]+]] "foo" // CHECK-DAG: [[foo_name:%[a-zA-Z0-9_]+]] = OpString "foo" // CHECK-DAG: [[data_name:%[a-zA-Z0-9_]+]] = OpString "data" // CHECK-DAG: [[void:%[a-zA-Z0-9_]+]] = OpTypeVoid // CHECK-DAG: [[uint:%[a-zA-Z0-9_]+]] = OpTypeInt 32 0 // CHECK: [[uint_0:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 0 -// CHECK-DAG: [[decl:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] Kernel [[foo]] [[foo_name]] -// CHECK-DAG: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] [[import]] ArgumentInfo [[data_name]] -// CHECK: OpExtInst [[void]] [[import]] ArgumentUniformTexelBuffer [[decl]] [[uint_0]] [[uint_0]] [[uint_0]] [[info]] +// CHECK-DAG: [[decl:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} Kernel [[foo]] [[foo_name]] +// CHECK-DAG: [[info:%[a-zA-Z0-9_]+]] = OpExtInst [[void]] {{.*}} ArgumentInfo [[data_name]] +// CHECK: OpExtInst [[void]] {{.*}} ArgumentUniformTexelBuffer [[decl]] [[uint_0]] [[uint_0]] [[uint_0]] [[info]]