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