Add ROCm support. #393
Add ROCm support. #393
Conversation
|
|
Codecov Report
@@ Coverage Diff @@
## master #393 +/- ##
==========================================
+ Coverage 62.42% 65.51% +3.09%
==========================================
Files 126 121 -5
Lines 7116 6540 -576
Branches 1255 1141 -114
==========================================
- Hits 4442 4285 -157
+ Misses 2469 2053 -416
+ Partials 205 202 -3
Flags with carried forward coverage won't be shown. Click here to find out more.
Continue to review full report at Codecov.
|
| __syncwarp(); | ||
| #endif | ||
| #ifdef __HIP_PLATFORM_HCC__ | ||
| // confused !!! |
There was a problem hiding this comment.
__syncwarp() in order to avoid read-write conflict of output_val, when warpReduceSum reading output_val from other threads within a warp.
AMD HIP doesn't support __syncwarp() now. Using __syncthreads() instead is ok, although bringing a few performance decrease.
| at::cuda::CUDAGuard device_guard(bboxes1.device()); | ||
| cudaStream_t stream = at::cuda::getCurrentCUDAStream(); | ||
| #endif | ||
| #ifdef __HIP_PLATFORM_HCC__ | ||
| // at::cuda::HIPGuard device_guard(bboxes1.device()); |
There was a problem hiding this comment.
Because pytorch handles ROCm as CUDA, it will cause assert device type error if not comment.
I'm not familiar with CUDA or ROCm, but i hope ROCM can get official efficient support.
There was a problem hiding this comment.
It seems that here should use HIPGuardImplMasqueradingAsCUDA.h and getCurrentHIPStreamMasqueradingAsCUDA?
|
closed by duplicate PR #1022 |
|
Hi @acai66！We are grateful for your efforts in helping improve this open-source project during your personal time. |
note:
I have modified some CUDA apis to HIP apis, so i can compile it with ROCm. I have test mmcv/tests/test_ops/test_bbox.py and test_nms.py, both of them work fine.
But i don't have all the tests, and i have only an AMD graphic card, so i can not test it with CUDA.
It needs to be tested enough Before merging to main stream.