[amd_hip_fp16.h] Remove anonymous namespace from __half functions #90
Conversation
danzimm
requested review from
chrispaquot,
gandryey,
saleelk,
mangupta and
rakesroy
as code owners
| @@ -530,1250 +530,1247 @@ THE SOFTWARE. | |||
| }; | |||
| // END STRUCT __HALF2 | |||
|
|
|||
| namespace | |||
There was a problem hiding this comment.
why do we need anonymous namespace here before? Is it for avoiding duplication with other definition?
There was a problem hiding this comment.
Really unsure— looking at the blame there isn’t any context to including the anon namespace unfortunately. Since the function parameters are different from other definitions with the same name we shouldn’t have any duplicate function name issues (thanks to mangling/function overloading)
|
I think we made some changes to fix the ODR violation in bfloat16 header. Can you please try that with latest HIP/ROCM. #include <Eigen/Core>
#include <amd_detail/amd_hip_fp16.h>
#include <amd_detail/amd_hip_bf16.h>
__global__ void kernel(float* res1, float* res2) {
auto hf1 = __float2half(2.718f);
auto bf1 = __float2bfloat16(2.718f);
*res1 = __half2float(hlog(hf1));
*res2 = __bfloat162float(hlog(bf1));
}It seems to be working fine. |
|
I reproduced this with ROCm 6.2 RC3-- did you have a different version in mind? Also just to confirm, what were your compiler flags for the above file? It was compiled with hipcc right? I think we repro'd with just clang. |
|
Can you share your minimum reproducer here. Regarding 6.2, It should have the fix, its just your original error, had the old hip_bf16 changes. The definition of hlog is: Regarding compilation flags, I just used |
|
Yes! Give me a few hours-- we got RC4 today so I'll use that to repro the issue. Thanks for the quick response here. |
|
This is my third major edit, apologies-- I had myself confused a few times with my repro. I now have a minimal repro with eigen 3.3.90 & rocm 6.2 rc4: and I get the following error: BTW I don't think I mentioned this from above-- I'm working on this for Meta |
|
Can you share your compile command as well. The code is as follows? #include <hip/hip_bf16.h>
#include <Eigen/Core>
int main() {} |
|
I think I can reproduce the issue, but having some trouble with my machine. I think I am missing some newer devicelibs. Reinstalling latest version of rocm and trying, will update it here. |
|
I think a better way is to fix this in eigen lib itself. hlog is the only function being referenced from global namespace. other half functions like hsqrt or hexp are referenced without any Raised https://gitlab.com/libeigen/eigen/-/merge_requests/1661/diffs to fix this on eigen, lets see what folks say on it. |
|
Sure, works for me (double checked with our repro). Thanks for putting up that patch. Going to go ahead and close this in hope that we can get that eigen patch upstreamed. |
When including both
amd_hip_fp16.handamd_hip_bf16.hin the same file and trying to use functions in the anonymous namespace insideamd_hip_fp16.hwhich are shadowed by functions in the global namespace ofamd_hip_bf16.h, e.g.hlog(__half)comes from the former and is shadowed byhlog(const __hip_bfloat16)from the latter, we're met with a compiler error of the form (the top file isHalf.h):This is due to the fact that the global definition hides the definition in the anonymous namespace. Looking at the git history of this file it doesn't appear there's any specific reason these functions need to be in an anonymous namespace so I went ahead and removed it so that e.g.
hlog(__half)participates in overload resolution.This patch works when using prebuilts from RPMs & patching the headers directly. I am unable to build from source to run tests locally. Let me what verification you want me to run: when I tried following the README or https://rocm.docs.amd.com/projects/HIP/en/latest/install/build.html to build & run tests I run into various errors that no docs refer to (it seems I need to clone a few other repos & possibly build another project first, but I can't quite figure out the right incantations-- I've gotten as far as a cmake error complaining about not finding
amd_comgr)