Skip to content

Fix translation of read_image* built-ins to SPIR-V #408

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
AlexeySotkin merged 1 commit into from
Jan 17, 2020
Merged

Fix translation of read_image* built-ins to SPIR-V #408

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
20 changes: 17 additions & 3 deletions lib/SPIRV/OCL20ToSPIRV.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Copy link
Member

@svenvh svenvh Jan 14, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Including the return type in the mangled name does not sound like the correct way to handle this issue.

SPIR-V 1.4 adds the SignExtend and ZeroExtend image operands. It would be better to make use of those somehow.

We have a downstream patch that uses a single bit in the image operands to disambiguate. I could try to upstream that instead if you are interested, but it will require a bit of rework and it will not be conformant to any pre-1.4 spec.

Copy link
Contributor

@AlexeySotkin AlexeySotkin Jan 15, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We also should disambiguate between the following 2 function calls:

uint4 c = read_imageui(input, (int4)(0, 0, 0, 0));
float4 f = read_imagef(input, (int4)(0, 0, 0, 0));

I don't think the image operands can help to represent these calls in SPIR-V-friendly LLVM IR.

Copy link
Member

@svenvh svenvh Jan 16, 2020

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ah sorry, I didn't realize this was only about the SPIR-V representation in LLVM IR.

LGTM then!

// 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;
Expand Down
24 changes: 24 additions & 0 deletions test/read_image.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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));
}