"Native API returns: -42 (PI_ERROR_INVALID_BINARY)" when running any example from joint_matrix extension

[Hitman4Reason]

When running examples of the joint_matrix extension(like sycl/test/matrix/matrix-int8-test.cpp) with "clang++ -fsycl -Xsycl-target-backend --cuda-gpu-arch=sm_75 -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 matrix-int8-test.cpp -o matri.out" it compiles but then I get
"terminate called after throwing an instance of 'sycl::_V1::runtime_error'
what(): Native API failed. Native API returns: -42 (PI_ERROR_INVALID_BINARY) -42 (PI_ERROR_INVALID_BINARY)
Aborted"
I have tried many combinations of the commands and flags to compile this for cuda but most return this error when running "./matri.out".

After building sycl with cuda backend I compile the "simple-sycl-app.cpp" with no issues with "clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda \ simple-sycl-app.cpp -o simple-sycl-app-cuda.exe" and run with "ONEAPI_DEVICE_SELECTOR=cuda:* ./simple-sycl-app-cuda.exe" but any example code that uses the joint_matrix extension I cannot get to run.

Environment

• OS: Linux
• Target device and vendor: Nvidia T4
• DPC++ version: clang version 18.0.0 (https://github.com/intel/llvm 2f90192)
• CUDA : Driver Version: 525.105.17 CUDA Version: 12.0

Thanks in advance for any help.

[JackAKirk]

If you look at the top of "sycl/test/matrix/matrix-int8-test.cpp" you'll see that this test is marked as not supported on ext_oneapi_cuda backend, so this behavior is expected.

The joint_matrix extension cuda support is fully tested by the following tests:

joint_matrix_tensorcores_sm70.cpp
joint_matrix_tensorcores_sm72.cpp
joint_matrix_tensorcores_sm80.cpp

Although obviously if you want to run on sm75 the sm80 test won't work.

With regard to the error, we are in the process of updating our troubleshooting section to explain this particular error code. For the time being this is the explanation:

Dealing with Invalid Binary Errors

Incorrect Platform

A common mistake is to execute a SYCL program using a platform for which
the SYCL program does not have a compatible binary. For example the SYCL
program may have been compiled for a SPIR-V backend but then executed on a
HIP device. In such a case the following error code,
PI_RESULT_ERROR_INVALID_BINARY, will be thrown. In this scenario, check
the following points:

1. Make sure your target platform is in -fsycl-targets so that the
   program will be compiled for the required platform(s).
2. Make sure that the program is using a sycl platform or device selector
   that is compatible with the platforms for which the executable was
   compiled.

[JackAKirk]

Actually I think there is a bug involved here: the test should not have been possible to compile on the cuda backend because it uses sizes of the joint_matrix that are invalid on that backend. I will investigate this. There should have been a compile time error stating that the configuration was invalid on that backend.

[JackAKirk]

Actually I think there is a bug involved here: the test should not have been possible to compile on the cuda backend because it uses sizes of the joint_matrix that are invalid on that backend. I will investigate this. There should have been a compile time error stating that the configuration was invalid on that backend.

Actually there is no bug in the impl. The problem arises with your compilation invocation which is not valid for the cuda backend: The above description of PI_RESULT_ERROR_INVALID_BINARY explains what is going on.
You used:

clang++ -fsycl -Xsycl-target-backend --cuda-gpu-arch=sm_75 -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 matrix-int8-test.cpp -o matri.out

where you missed the -fsycl-targets=nvptx64-nvidia-cuda, A correct instruction adds this to your invocation:

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_75 -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 matrix-int8-test.cpp -o matri.out

Which then leads to the compile time error as I wrote about above in the previous message.

If you have found our release docs/gettingStartedGuide.md documentation unhelpful then please let us know your feedback so we can improve it. Thanks. If you failed to find this documentation then your feedback is also welcome.

[JackAKirk]

Also, all the tests in "sycl/test" are not end-to-end tests for the sycl runtime. They are only testing device code generation and are inappropriate for execution. End to end tests appropriate for execution are all located in "sycl/test-e2e"

[Hitman4Reason]

Also, all the tests in "sycl/test" are not end-to-end tests for the sycl runtime. They are only testing device code generation and are inappropriate for execution. End to end tests appropriate for execution are all located in "sycl/test-e2e"

Thanks for all the clarifications. I completelly missed the fact that those examples were not meant for execution on cuda. I had used the compiler flags you suggested aswell but since at that point it didn't even compile I thought that it wasn't the right way either. Now I think everything is clarified.

[Hitman4Reason]

Just as a pointer also, I think there is a typo in https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_oneapi_matrix.asciidoc#example-using-int8_t-type since the joint_matrix_mad uses 5 arguments instead of 4 but I might be mistaken.

[JackAKirk]

Just as a pointer also, I think there is a typo in https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_matrix/sycl_ext_oneapi_matrix.asciidoc#example-using-int8_t-type since the joint_matrix_mad uses 5 arguments instead of 4 but I might be mistaken.

Well spotted, but this is because the interface was recently changed in the specification: #11004
This was to support the cases where the type of the data in the output matrix differs from the Accumulator matrix. This is therefore not really a typo, it is just that the specification is ahead of the implementation, that I believe has not yet been changed.

[JackAKirk]

Spec is not compliant with implementation.