From 10a4c4e74ef87e30034b6f3d1b30935ffb809169 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 14 Jan 2020 16:12:36 +0300 Subject: [PATCH] Fix translation of read_image* built-ins to SPIR-V 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 --- lib/SPIRV/OCL20ToSPIRV.cpp | 20 +++++++++++++++++--- test/read_image.cl | 24 ++++++++++++++++++++++++ 2 files changed, 41 insertions(+), 3 deletions(-) create mode 100644 test/read_image.cl diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp index 75a6b58d22..3c07a8c4a2 100644 --- a/lib/SPIRV/OCL20ToSPIRV.cpp +++ b/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; diff --git a/test/read_image.cl b/test/read_image.cl new file mode 100644 index 0000000000..48715e93f6 --- /dev/null +++ b/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)); +}