Backport fix translation of read_image* built-ins to SPIR-V to 9.0 re…

…lease

There are several overloads of read_image function defined by OpenCL C
spec, but all of them use the same SPIR-V instruction, so, we need to
add one more optional postfix to differentiate instructions with the
same argument types, but different return types.

Example:
- int4 read_imagei(image2d_t, sampler_t, int2)
- float4 read_imagef(image2d_t, sampler_t, int2)

Both functions above are represented by the same SPIR-V instruction and
we need to distinguish them in SPIR-V friendly LLVM IR.

Signed-off-by: Alexey Sachkov <[email protected]>

lib/SPIRV/OCL20ToSPIRV.cpp

@@ -1027,9 +1027,23 @@ void OCL20ToSPIRV::transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
   unsigned ExtOp = ~0U;
   if (StringRef(Info.UniqName).startswith(kSPIRVName::Prefix))
     return;
-  if (OCLSPIRVBuiltinMap::find(Info.UniqName, &OC))
-    Info.UniqName = getSPIRVFuncName(OC);
-  else if ((ExtOp = getExtOp(Info.MangledName, Info.UniqName)) != ~0U)
+  if (OCLSPIRVBuiltinMap::find(Info.UniqName, &OC)) {
+    if (OC == OpImageRead) {
+      // There are several read_image* functions defined by OpenCL C spec, but
+      // all of them use the same SPIR-V Instruction - some of them might only
+      // differ by return type, so, we need to include return type into the
+      // mangling scheme to get them differentiated.
+      //
+      // Example: int4 read_imagei(image2d_t, sampler_t, int2)
+      //          uint4 read_imageui(image2d_t, sampler_t, int2)
+      // Both functions above are represented by the same SPIR-V
+      // instruction: argument types are the same, only return type is
+      // different
+      Info.UniqName = getSPIRVFuncName(OC, CI->getType());
+    } else {
+      Info.UniqName = getSPIRVFuncName(OC);
+    }
+  } else if ((ExtOp = getExtOp(Info.MangledName, Info.UniqName)) != ~0U)
     Info.UniqName = getSPIRVExtFuncName(SPIRVEIS_OpenCL, ExtOp);
   else
     return;

test/read_image.cl

@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -triple spir64 -finclude-default-header -O0 -cl-std=CL2.0 -emit-llvm-bc %s -o %t.bc
+// RUN: llvm-spirv %t.bc -o %t.spv
+// RUN: spirv-val %t.spv
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv -s %t.bc -o %t1.bc
+// RUN: llvm-dis %t1.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
+
+// CHECK-SPIRV: TypeInt [[IntTy:[0-9]+]]
+// CHECK-SPIRV: TypeVector [[IVecTy:[0-9]+]] [[IntTy]]
+// CHECK-SPIRV: TypeFloat [[FloatTy:[0-9]+]]
+// CHECK-SPIRV: TypeVector [[FVecTy:[0-9]+]] [[FloatTy]]
+// CHECK-SPIRV: ImageRead [[IVecTy]]
+// CHECK-SPIRV: ImageRead [[FVecTy]]
+
+// CHECK-LLVM: call spir_func <4 x i32> @_Z24__spirv_ImageRead_Ruint414ocl_image3d_roDv4_i
+// CHECK-LLVM: call spir_func <4 x float> @_Z25__spirv_ImageRead_Rfloat414ocl_image3d_roDv4_i
+
+__kernel void kernelA(__read_only image3d_t input) {
+  uint4 c = read_imageui(input, (int4)(0, 0, 0, 0));
+}
+
+__kernel void kernelB(__read_only image3d_t input) {
+  float4 f = read_imagef(input, (int4)(0, 0, 0, 0));
+}