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..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"; @@ -165,6 +169,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 +202,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..b1390a394 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,6 +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. + // 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; @@ -2516,6 +2521,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 +2554,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; @@ -7164,6 +7173,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 +7195,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/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 new file mode 100644 index 000000000..3efb908b2 --- /dev/null +++ b/test/Reflection/readwrite_texel_buffer_argument.cl @@ -0,0 +1,16 @@ +// 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: 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]] {{.*}} 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 new file mode 100644 index 000000000..134b7bc02 --- /dev/null +++ b/test/Reflection/storage_texel_buffer_argument.cl @@ -0,0 +1,17 @@ +// 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 + +kernel void foo(write_only image1d_buffer_t data) { +} + +// 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]] {{.*}} 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 new file mode 100644 index 000000000..46de30453 --- /dev/null +++ b/test/Reflection/uniform_texel_buffer_argument.cl @@ -0,0 +1,18 @@ +// 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 + +kernel void foo(read_only image1d_buffer_t data) { +} + +// 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]] {{.*}} 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]] +