Errors during compilation of joint_matrix_tensorcores.cpp

[vidyalatha-badde]

Hi,

I'm trying to compile the sample code provided here https://github.com/intel/llvm-test-suite/blob/intel/SYCL/Matrix/joint_matrix_tensorcores.cpp with cuda backend with the following compilation command (please correct me if there is any issue with the compilation command)
clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_80 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 joint_matrix_tensorcore.cpp

I'm getting a lot of errors during the compilation. Please see the error log attached.
error_log.txt
Could someone help me out in compiling the code successfully?

Thanks in advance!

~Vidya

[JackAKirk]

Hi,

This is because you are using the 2023.0 release. The joint_matrix extension is not officially supported in this release for the cuda backend. If you wish to use joint_matrix with the cuda backend using 2023.0 then you need to use the "legacy" joint_matrix cuda interface that it different from the portable interface that can be used with both Intel and cuda backends which will be available in 2023.1: an example of the 2023.0 version of the joint_matrix cuda backend APIs and details for how to compile (see the top of the file) it are given in this test case: https://github.com/intel/llvm-test-suite/blob/intel/SYCL/Matrix/joint_matrix_tensorcores_legacy.cpp

[vidyalatha-badde]

Hello @JackAKirk ,

Many thanks for your reply!
I just tried compiling the legacy code which you pointed out but still i get the errors (attaching the log here)
Compilation command used:
clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_80 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=3 legacy.cpp
error_log_legacy.txt
Since the errors are related to use of undeclared identifier 'bfloat16' bfloat16>::value)
tried adding the following lines to the code

#include <ext/oneapi/experimental/bfloat16.hpp>
#include <ext/oneapi/matrix/matrix.hpp>
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;

now the undeclared errors are gone but got some new errors
error_log_with_modifications.txt

Could you please help me with the compilation of the code successfully without any errors?
Also, it would be a great help if you could let me know on which hardware (both Intel and nvidia) can we use this joint_matrix API. (how to check if a particular hardware is compatible or not to work with joint_matrix API?)

Once again thanks for your time.

~Vidya

[AlexeySachkov]

Tagging @dkhaldi here to help reporter with joint_matrix APIs

[vidyalatha-badde]

Hello @dkhaldi ,

Could you please take a look at the issue and help me here?

~Vidya.

[JackAKirk]

Hello @JackAKirk ,

Many thanks for your reply! I just tried compiling the legacy code which you pointed out but still i get the errors (attaching the log here) Compilation command used: clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_80 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=3 legacy.cpp error_log_legacy.txt Since the errors are related to use of undeclared identifier 'bfloat16' bfloat16>::value) tried adding the following lines to the code

#include <ext/oneapi/experimental/bfloat16.hpp> #include <ext/oneapi/matrix/matrix.hpp> using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;

now the undeclared errors are gone but got some new errors error_log_with_modifications.txt

Could you please help me with the compilation of the code successfully without any errors? Also, it would be a great help if you could let me know on which hardware (both Intel and nvidia) can we use this joint_matrix API. (how to check if a particular hardware is compatible or not to work with joint_matrix API?)

Once again thanks for your time.

~Vidya

Hi, I think the errors are due to const T cases not being supported yet in 2023.0. You can try to comment out these cases in main. The full list of types/sizes supported in the cuda backend is given here: https://github.com/intel/llvm/blob/e50a2f5f97acb12db1de78c9ad739b931c77b03f/sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix/sycl_ext_oneapi_cuda_matrix.asciidoc
All supported types are tested in joint_matrix_tensorcores.cpp.

[vidyalatha-badde]

Hello @JackAKirk

Thanks!

Yes, you are right I commented out the test cases with const T and now I don't see any compilation errors in the legacy code.
Running the executable doesn't give any output, so I tried printing the computeCapability for my device, and below is the output

./a.out
computeCapability: 6

PI CUDA ERROR:
        Value:           209
        Name:            CUDA_ERROR_NO_BINARY_FOR_GPU
        Description:     no kernel image is available for execution on the device
        Function:        build_program
        Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:709


PI CUDA ERROR:
        Value:           400
        Name:            CUDA_ERROR_INVALID_HANDLE
        Description:     invalid resource handle
        Function:        cuda_piProgramRelease
        Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:3558

terminate called after throwing an instance of 'sycl::_V1::compile_program_error'
  what():  The program was built for 1 devices
Build program log for 'Tesla P100-PCIE-12GB':
 -999 (Unknown PI error)
Aborted (core dumped)

Since the device compute capability is 6.0 I think the code doesn't work for my target device i.e. Tesla P100-PCIE-12GB.
Am I correct?

What is the minimum compute capability required to make joint_matrix API's work on Nvidia GPU and are there any such known limitations to make it work on Intel hardware as well?

It would be great if you could help me out with the above information.

Regards,
Vidya.

[JackAKirk]

Hi,

You can see the table giving the compute compatibility for each case in https://github.com/intel/llvm/blob/e50a2f5f97acb12db1de78c9ad739b931c77b03f/sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix/sycl_ext_oneapi_cuda_matrix.asciidoc

You can also see this from the joint_matrix_tensorcores.cpp example: each case is only running if the check for compute capability passes the minimum for each operation (e.g. https://github.com/intel/llvm-test-suite/blob/intel/SYCL/Matrix/joint_matrix_tensorcores_legacy.cpp#L238): the absolute minimum is sm_70, which is why you won't be able to run any of the cases on a sm_6x card like a P100.

Best wishes

[vidyalatha-badde]

Hi @JackAKirk,

I just checked with the latest version of oneAPI 2023.1 and I could successfully compile the joint_matrix_tensorcores.cpp example code without any errors.
Since I do not have a cuda device whose compute capability is > 7, can I run the code on Intel devices?
Will this code work for Intel CPUs/GPUs (if yes could you please let me know the compilation command to be used)?

Kindly let me know.

Regards,
Vidya.

[vidyalatha-badde]

Hi @JackAKirk,

A quick update:
Finally, I'm able to find a CUDA device whose compute capability is 7.5 (Tesla T4). when I try to run the joint_matrix_tensorcores.cpp sample code the execution fails after it enters the if (computeCapability >= 7.0) { statement.
Here is the screenshot of the output

and gdb output
gdb_output.txt

Could you please check this sample code and let me know what is missing from my end?

Regards,
Vidya.

[JackAKirk]

Could you please check this sample code and let me know what is missing from my end?

Regards, Vidya.

Hello,

Please see the information that I have added here: #9019
It explains why you saw the above error. The latest up to date tests for joint_matrix with the ext_oneapi_cuda backend are now here:
https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/Matrix/joint_matrix_tensorcores_sm70.cpp
https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/Matrix/joint_matrix_tensorcores_sm72.cpp
https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/Matrix/joint_matrix_tensorcores_sm80.cpp

Each of these three tests runs all the test cases that correspond to the supported parameters combinations (that I detail in the table added in #9019) that require a minimum compute capability of sm_70, sm_72 and sm_80 respectively.
Since you are running on sm_75, this means that you will be able to compile and execute the tests using the Tensor Cores contained within the joint_matrix_tensorcores_sm70.cpp and joint_matrix_tensorcores_sm72.cpp cases. For the joint_matrix_tensorcores_sm80.cpp you will be able to compile the test, however the test cases will be skipped.
As detailed in #9019, for your sm_75 case you will be able to compile joint_matrix_tensorcores_sm70.cpp and execute if you provide the additional flag -Xsycl-target-backend --cuda-gpu-arch=sm_xx where sm_xx is either sm_70, sm_72 or sm_75.
Similarly, you will be able to compile joint_matrix_tensorcores_sm72.cpp and execute if you provide the additional flag -Xsycl-target-backend --cuda-gpu-arch=sm_xx where sm_xx is either sm_72 or sm_75.

If you would like to know which parameter combinations are supported on Intel XMX or AMX devices consult the corresponding tables added to the same document in this PR: #7964.