[SYCL] Enable AMD GPU support. #3795
Conversation
malixian
requested review from
AaronBallman,
AGindinson,
bader,
elizabethandrews,
mdtoguchi,
premanandrao,
smaslov-intel and
a team
as code owners
| @@ -141,6 +150,7 @@ def main(): | |||
| parser.add_argument("-t", "--build-type", | |||
| metavar="BUILD_TYPE", default="Release", help="build type: Debug, Release") | |||
| parser.add_argument("--cuda", action='store_true', help="switch from OpenCL to CUDA") | |||
| parser.add_argument("--rocm", action='store_true', help="swith from OpenCL to ROCM") | |||
There was a problem hiding this comment.
@malixian, one note.
As of today, we don't have AMD GPU HW in our CI system, so we won't be able to validate that other contributions do not break ROCM support.
I think we should decide ASAP, how ROCM support is going to be verified.
There was a problem hiding this comment.
This is indeed a problem.
There was a problem hiding this comment.
At least for plugin we can use hip<->cuda source code compatibility
|
@malixian, thanks for contribution! Please update documentation, https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md at least. |
All right. |
|
We update the code and GetStartedGuide.md document for AMDGPU. |
There was a problem hiding this comment.
Per discussion with @bader, adding & committing 3 suggestions (tested locally) to:
- Apply the Driver LIT patch correctly to pass the checks
- Remove
__SYCL_AMDGCN__macro due to the similar efforts for__SYCL_NVPTX__in [SYCL] Remove the __SYCL_NVPTX__ macro #3977
Fixed comments.
|
If |
|
Enable AMG GPU for DPC++. To support this feature, We did two parts of development, namely the compilation tool chain and the runtime. We have implemented rocm-plugin refer to cuda-plugin. Many test cases in this project have passed, but there are still some problems which can be roughly divided into three areas:
First, some errors occurred at link time. like all-pairs-disance-sycl case occurred error:
undefined hidden symbol: __spirv_ControlBarrier will appear during the lld link perioddue to usebarrier(access::fence_space::local_space). Similarly, theundefined__spirv_SubgroupShuffleINTELerror occurs when thecl::sycl::atomickeyword is used at that time.Second, some errors occurred at runime. the program will core dump when calling
hipMemcpyDtoHAsyncAPI due to allocating memory size is too large. But cuda does not have this problem.Finally, calculation accuracy problem. Currently, the calculation of
floattype kernel functions is inaccurate, and there is no problem withinttype testing. We haven't figured out where the problem is.We will keep track of the above issues.