-
Notifications
You must be signed in to change notification settings - Fork 463
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[ROCm] bfloat16 enablement #1560
Conversation
✅ Deploy Preview for pytorch-fbgemm-docs canceled.
|
LGTM; importing to Phabricator |
@q10 has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
@q10 has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
@liligwu Sorry to bug you one more time. Could I trouble you to do 2 things to keep our build system happe:
@@ -1398,12 +1398,13 @@
};
#ifdef __HIP_PLATFORM_HCC__
-using __nv_bfloat16=hip_bfloat16;
+using __nv_bfloat16 = hip_bfloat16;
typedef struct __align__(4) {
uint16_t x;
uint16_t y;
-} __nv_bfloat162_raw;
+}
+__nv_bfloat162_raw;
struct __align__(4) __nv_bfloat162 {
__nv_bfloat16 x;
@@ -1479,12 +1480,12 @@
#ifdef __HIP_PLATFORM_HCC__
// the descriptions of __float2bfloat16 and __float2bfloat16_rn are identical
// https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH____BFLOAT16__MISC.html#group__CUDA__MATH____BFLOAT16__MISC
-static __host__ __device__ __nv_bfloat16 __float2bfloat16(float f){
+static __host__ __device__ __nv_bfloat16 __float2bfloat16(float f) {
__nv_bfloat16 output;
return output.round_to_bfloat16(f);
}
-static __host__ __device__ __nv_bfloat16 __float2bfloat16_rn(float f){
+static __host__ __device__ __nv_bfloat16 __float2bfloat16_rn(float f) {
__nv_bfloat16 output;
return output.round_to_bfloat16(f);
} After this, we should be good to merge. Thanks for your patience! |
@q10 , it's done. |
@q10 , @shintaro-iwasaki . Could you please close this PR if everything looks good? Thank you. |
@q10 has imported this pull request. If you are a Meta employee, you can view this diff on Phabricator. |
Yep, was making sure the GitHub builds were passing |
Regarding the ROCm CI failure
I opened an issue #1559 I don't know if it was observed in CI, but it failed in a CUDA container I tested. |
Thanks for filing the issue. The issue is observed in CI, but it is very unreliable, as one of the earlier builds in this PR had previously passed prior to your last lint commit. I suspect it is an issue with the instance that the job was running on. We will investigate this further after this PR is merged. |
Enable bfloat16 on ROCm