[PATCH] D142907: LangRef: Add "dynamic" option to "denormal-fp-math"

[Matt Arsenault via Phabricator via cfe-commits]

arsenm marked 2 inline comments as done.
arsenm added inline comments.


================
Comment at: clang/lib/CodeGen/CGCall.cpp:2059-2061
+  F.removeFnAttrs(AttrsToRemove);
+  addDenormalModeAttrs(Merged, MergedF32, FuncAttrs);
+  F.addFnAttrs(FuncAttrs);
----------------
tra wrote:
> IIUIC, this changes denorm mode attributes on the functions with dynamic denorm mode that we link in.
> Will that be a problem if the same function, when linked into different modules, would end up with different attributes? E.g. if a function is externally visible and is intended to be common'ed across multiple modules. Should dynamic denorm mode be restricted to the functions private to the module only? We do typically internalize linked bitcode for CUDA, but I don't think it's something we can always implicitly assume.
> 
The whole point of -mlink-builtin-bitcode is to apply the attributes for the current compilation to what's linked in. The linked functions are always internalized. The only case where we might not want to internalize is for weak symbols (but it looks like we do internalize those today, but this is something I've thought about changing). I'll add a test with a weak library function

In the weak case the right thing to do is probably to not change from dynamic, simply because this linking process is outside of the user's control.




================
Comment at: clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu:15
+// RUN:   %S/Inputs/ocml-sample.cl -o %t.dynamic.full.bc
+
+
----------------
tra wrote:
> Do we want to verify that the compiled samples have the correct function attributes?
Maybe, but that's already tested separately. This test is a bit complex as it is (and could maybe use a few more combinations)


================
Comment at: clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu:77
+
+// CHECK: kernel_f32({{.*}}) #[[$KERNELATTR:[0-9]+]]
+__global__ void kernel_f32(float* out, float* a, float* b, float* c) {
----------------
tra wrote:
> Nit: CHECK-LABEL ?
 error: found 'CHECK-LABEL:' with variable definition or use



================
Comment at: clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu:94
+// We should not be littering call sites with the attribute
+// CHECK-NOT: denormal-fp-math
+
----------------
tra wrote:
> I'm not sure whether it does what it's intended to. 
> AFAICT, at this point we will be past the call sites, so if it's intended to check the call sites in kernel_*, it will likely always succeed, even if we do litter call sites with unwanted attributes.
> 
> It's also possible that I have a wrong idea about what the expected IR looks like. If you could post it for reference, that would be helpful.
I can drop this, I later added the -implicit-check-not=denormal-fp-math to all the FileChecks


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D142907/new/

https://reviews.llvm.org/D142907