[SYCL] Improve the error mechanism of llvm-no-spir-kernel

[AGindinson]

This patch improves the tool's diagnostic upon finding a
SPIR kernel within an LLVM module. Despite that the tool's
only current use is within the SYCL FPGA flow, it's important
to make the message target-agnostic, so that the tool is not
tied to a particular device BE.
A related commit to the Clang driver has extended these diagnostics
with SYCL FPGA specifics without affecting the tool itself.

This patch also introduces testing for the return code value. For
example, this should allow the Clang driver users/developers to
differentiate between the two possible causes of llvm-no-spir-kernel
failure.

Signed-off-by: Artem Gindinson [email protected]

[Fznamznon]

I do not have objections against the patch. But I do not see purpose of these changes (as well as the purpose of the tool). The commit message is confusing:

This patch improves the tool's diagnostic upon finding a
SPIR kernel within an LLVM module. Despite that the tool's
only current use is within the SYCL FPGA flow, it's important
to make the message target-agnostic - this way, the tool retains
the potential to be re-used for other targets with an LLVM-based
backend.

What exactly do you mean here? If other targets is other programming models I do not see how it can be reused for other targets, it checks SPIR_KERNEL calling convention, AFAIK it appears only in spir-based IR targets which is used only by OpenCL/SYCL. But if under other targets you mean non-FPGA targets SYCL, I don't see why it cannot be used for them now.

[AGindinson]

But I do not see purpose of these changes (as well as the purpose of the tool).

This tool is called by the Clang driver in the FPGA flow. Before an object file is linked against an archive that contains AOT-compiled device code, the tool can detects the presence of a SPIRKernel-calling convention function in the object file. This would mean that the device code from the object file was not AOT-compiled, neither stored in the archive. Terminating the compilation in such case and requesting the archive re-compilation to include all device code is better than letting the user run into a runtime error due to an "unknown kernel".

The purpose of these changes is to provide a helpful diagnostic message to the user - from within the tool, we can reach the name of the missing function, and include it in the diagnostic. The other "half" of the diagnostic will be provided by the Clang driver, giving a more detailed explanation with FPGA specifics. I will be uploading the driver patch shortly.

But if under other targets you mean non-FPGA targets SYCL,

Yes, non-FPGA SYCL targets are implied.

I don't see why it cannot be used for them now.

It can. The point was "let's not affect the tool's re-usability for non-FPGA targets by adding FPGA-specific diagnostics".

I'll reformulate the commit message into something less confusing.

[AGindinson]

Updated the commit message & the PR description.

[Fznamznon]

This tool is called by the Clang driver in the FPGA flow. Before an object file is linked against an archive that contains AOT-compiled device code, the tool can detects the presence of a SPIRKernel-calling convention function in the object file. This would mean that the device code from the object file was not AOT-compiled, neither stored in the archive. Terminating the compilation in such case and requesting the archive re-compilation to include all device code is better than letting the user run into a runtime error due to an "unknown kernel".

I'd say that I do not see why we do that. This relies on presence of only one module with device code. But we could support cases when there are multiple device modules.

The purpose of these changes is to provide a helpful diagnostic message to the user - from within the tool, we can reach the name of the missing function, and include it in the diagnostic. The other "half" of the diagnostic will be provided by the Clang driver, giving a more detailed explanation with FPGA specifics. I will be uploading the driver patch shortly.

Are you sure that Unexpected SPIR kernel occurence: typeinfo name for run_kernel_add(cl::sycl::queue&, float, float, float&)::$_0::operator()(cl::sycl::h andler&) const::SimpleAdd is some sort of useful diagnostic?

I'll reformulate the commit message into something less confusing.

It feels a bit better now.

[AGindinson]

Are you sure that Unexpected SPIR kernel occurence: typeinfo name for run_kernel_add(cl::sycl::queue&, float, float, float&)::$_0::operator()(cl::sycl::h andler&) const::SimpleAdd is some sort of useful diagnostic?

Combined with the intended Clang driver changes, it should make sense. I've decided to split the changes into the two patches as per the "one commit - one component" policy. I'll provide the link to the driver PR ASAP.

[AGindinson]

This tool is called by the Clang driver in the FPGA flow. Before an object file is linked against an archive that contains AOT-compiled device code, the tool can detects the presence of a SPIRKernel-calling convention function in the object file. This would mean that the device code from the object file was not AOT-compiled, neither stored in the archive. Terminating the compilation in such case and requesting the archive re-compilation to include all device code is better than letting the user run into a runtime error due to an "unknown kernel".

I'd say that I do not see why we do that. This relies on presence of only one module with device code. But we could support cases when there are multiple device modules.

In the SYCL FPGA flow, a separate llvm-no-spir-kernel call is made for each device module-containing object.

[Fznamznon]

This tool is called by the Clang driver in the FPGA flow. Before an object file is linked against an archive that contains AOT-compiled device code, the tool can detects the presence of a SPIRKernel-calling convention function in the object file. This would mean that the device code from the object file was not AOT-compiled, neither stored in the archive. Terminating the compilation in such case and requesting the archive re-compilation to include all device code is better than letting the user run into a runtime error due to an "unknown kernel".

I'd say that I do not see why we do that. This relies on presence of only one module with device code. But we could support cases when there are multiple device modules.

In the SYCL FPGA flow, a separate llvm-no-spir-kernel call is made for each device module-containing object.

I think there is some misunderstanding. Why do you think user can see a runtime error due to an "unknown kernel", if not all modules are stored in archive?

[AGindinson]

In the SYCL FPGA flow, a separate llvm-no-spir-kernel call is made for each device module-containing object.

I think there is some misunderstanding. Why do you think user can see a runtime error due to an "unknown kernel", if not all modules are stored in archive?

My earlier comment:

This would mean that the device code from the object file was not AOT-compiled, neither stored in the archive.

Since JIT is not supported for the targeted FPGA devices, it would not be possible to JIT a kernel from an outside object file. Any request to execute such a kernel from the in-archive host code would lead to the "unknown kernel" error.

[AGindinson]

I'll provide the link to the driver PR ASAP.

#1077

[AGindinson]

Squashed & updated the main commit message.

[mlychkov]

LGTM

[mdtoguchi]

LGTM

[AGindinson]

Note: the Linux failure is an unrelated sycl/test/assert.cpp one.

[bader]

Note: the Linux failure is an unrelated sycl/test/assert.cpp one.

@asavonic, IIRC, you were going to fix this. Any updates?