Differential D61399
[OpenMP][Clang] Support for target math functions Authored by gtbercea on May 1 2019, 2:01 PM.
Details
In this patch we propose a temporary solution to resolving math functions for the NVPTX toolchain, temporary until OpenMP variant is supported by Clang. We intercept the inclusion of math.h and cmath headers and if we are in the OpenMP-NVPTX case, we re-use CUDA's math function resolution mechanism. Authors:
Diff Detail
Event TimelineComment Actions For the record, this is an implementation of the scheme proposed in https://reviews.llvm.org/D60907#1484756. Finally, there is a remote chance this will cause trouble to people that use math.h/cmath functions, e.g. with the old SSE hack, which are not available anymore.
Comment Actions I don't like this implementation. Seems to me, it breaks one of the OpenMP standard requirements: the program can be compiled without openmp support. I assume, that with this includes the program won't be able to be compiled without OpenMP support anymore because it may use some device-specific math functions explicitly. Comment Actions Moreover, I think this will cause troubles even in simple cases. Assume we have target if(cond) construct. In this case we will need to compile the target region for both, the device and the host. If the target region uses some device-specific math functions, it will break the compilation for the host. Comment Actions Can you provide an example of a conforming program that can't be compiled without OpenMP support? Regardless of the use of any device-specific functions (which isn't covered by the standard, of course, but might be needed in practice), the code still needs to be compilable by the host in order to generate the host-fallback version. This doesn't change that. Thus, any program that uses anything from this math.h, etc. needs to compile for the host, and thus, likely compiles without OpenMP support. Maybe I'm missing your point, however. Comment Actions Assume we have something like this: #pragma omp target if(cond) a = __nv_xxxx(....); Instead of __nv_xxx you can try to use any Cuda-specific function, which is not the part of the standard math.h/cmath files. Will it be compilable even with OpenMP? Comment Actions I don't think that this changes that one way or the other. Your example won't work, AFAIK, unless you do something like: #pragma omp target if(cond) #ifdef __NVPTX__ a = __nv_xxxx(....); #else a = something_on_the_host; #endif and anything from these headers that doesn't also have a host version will suffer the same fate: if it won't also compile for the host (one way or another), then it won't work. Comment Actions The problem with this header file is that it allows to use those Cuda-specific functions unconditionally in some cases: #pragma omp target a = __nv_xxxx(....); It won't require any target-specific guards to compile this code (if we compile it only for Cuda-specific devices) and we're loosing the consistency here: in some cases target regions will require special device guards, in others, with the same function calls, it is not. And the worst thing, is that we implicitly allow to introduce this kind of incostistency into users code. That's why I would prefer to see a special kind of the include file, NVPTX specific, that must be included explicitly, so the user explictly commanded to use some target-specific math functions, if he really wants it. Plus, maybe, in this files we need force check of the platform and warn users that the functions from this header file must be used using device-specific checks. Or provide some kind of the default implementations for all the platforms, that do not support those math function natively. Comment Actions I believe that I understand your point, but two things:
#ifdef __NVPTX__
int __nv_floor();
#endif
int main() {
#pragma omp target
__nv_floor();
}and try to compile using Clang with -fopenmp -fopenmp-targets=nvptx64, the compilation fails: int1.cpp:8:1: error: use of undeclared identifier '__nv_floor' and this is because, when we invoke the compilation for the host, there is no declaration for that function. This is true even though nvptx64 is the only target for which the code is being compiled (because we always also compile the host fallback).
I think that you're getting at what I mentioned earlier, where if you have code that uses some function called, for example, rnorm3d -- this is a non-standard function, and so a user is free to define it, and then they compile with OpenMP target and include math.h, then the version of rnorm3d we bring in from the CUDA header will silently override their version (which would otherwise be implicitly declare target). But I don't think that this will happen either with this patch, instead, they'll get a warning about conflicting definitions (error: redefinition of 'rnorm3d'), and if it's an external-linkage function, then something should give them an error (although we should check this). The more interesting case, I think, actually comes when we switch to using declare variant, because then I think this silent override does occur, so we'd want to add a warning for when a variant delcared in a system header file would silently override a plain (non-variant) function declare/defined elsewhere. I believe that will give us all of the benefits of this while also addressing the concern you highlight. Comment Actions Ahh, yes, I forgot that we still generate the host version of the target region. Still, I think we need to prvide the default implementation of those non-standard functions (they can be very simple, maybe reporting error is going to be enough), which can be overriden by user. Comment Actions Did you test it for the builtins? like pow, powf and powl? How are the builtins resolved with this patch? Comment Actions
I appreciate your motivation, and I agree with you to some extent. I don't object to having generic versions of useful math functions, but I don't think they should be required. It's not reasonable to make someone add generic versions of every function which happens to appear in a system/target-specific math.h header. NVPTX won't be the only target that has target-optimized functions that get pulled in, even from our own headers, but system headers also have differences anyway depending on what preprocessor macros are defined. In the end, people can write portable code if they stick to what's in the standard, and we should make it reasonably easy for them to step outside of the standard to do what they need to do when the standard subset of available functionality doesn't meet their needs for whatever reason. This is what we do for C/C++, where we provide intrinsics and other system functions for those who can't write their code only using the facilities that C/C++ provide. In any case, I think that we can figure out how to add generic versions of non-standard math functions in a separate thread. I think that we should move forward with this and then make progress on generic versions separately. It's also possible that we want to fold this discussion into the discussion on an LLVM math library (we've talked about this for some time in the context of vector math libraries, and I'd not thought about accelerators in this context, but maybe this is all related). Comment Actions It is up to you. I don't have strong objections if you think this will work as required. Just the tests must be fixed, especially codegen tests.
Comment Actions
Thanks, Alexey. I think this will work as required, and then we'll be able to update it when we get declare variant. Agreed on the tests (on all points).
Comment Actions I've reverted this in rL360192 because it breaks stage 2 builds on GreenDragon. Please see the commit message and the inline comment for more details.
| |||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||
Please add a driver test covering this.