implement image1d_buffer using texel buffer

include/clspv/ArgKind.h

@@ -31,6 +31,8 @@ enum class ArgKind : int {
  Sampler,
  PointerUBO,
  PointerPushConstant,
+ StorageTexelBuffer,
+ UniformTexelBuffer,
 };
 
 // Converts an ArgKind to its string name.

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

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;

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<kTypes>(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;

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

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]]

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]]

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]]
+