[PATCH] D40573: [NVPTX] Assign valid global names

Artem Belevich via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Thu Nov 30 11:56:36 PST 2017


tra added inline comments.


================
Comment at: lib/IR/ValueSymbolTable.cpp:54-59
+      // MSVC name mangling but is fine for PTX.
+      if (M && Triple(M->getTargetTriple()).isNVPTX())
+        S << "$";
+      else
+        S << ".";
+    }
----------------
Hahnfeld wrote:
> tra wrote:
> > Hahnfeld wrote:
> > > tra wrote:
> > > > Hahnfeld wrote:
> > > > > tra wrote:
> > > > > > Hahnfeld wrote:
> > > > > > > tra wrote:
> > > > > > > > This patch addresses "we can't compile generated PTX because LLVM uses illegal characters", but exposes another issue -- having potentially different names on host and device is a problem for CUDA. For some objects host side may need to know what it's called on device side. We need it in order to access it from host (eg cudaMemcpyToSymbol(), or initializing static variables) and we currently assume that the names are the same. If such symbol gets different names on host and device, compilation will succeed, but we'll have problems at runtime.
> > > > > > > > 
> > > > > > > > Does "." have any special meaning? Can we skip the unique delimiter altogether? 
> > > > > > > > 
> > > > > > > > If we can't find a suitable way to guarantee identical naming, we'll need a way to have a reliable way to determine the name used on the other side of the compilation.
> > > > > > > > 
> > > > > > > > 
> > > > > > > So the interesting question is: When will this code ever be hit? Most programming languages (including C and C++) obviously don't allow multiple variables of the same name - how would the compiler say which symbol you meant. The use case I've mostly seen is for compiler generated function, `omp_outlined` for example. These can be generated multiple times in the same translation unit and have to get unique names. Do you have another example where this could happen?
> > > > > > > 
> > > > > > > I'm not really sure '.' has a special meaning. Maybe @rafael can help because one of his old commits (https://reviews.llvm.org/rL253804) says `For globals it is important to use "foo.1" to help C++ name demangling.`
> > > > > > I vaguely recall that '.' was an indication for demangler that it should not proceed further. I.e. a sort-of-special character to indicate the end of the C++-mangled part of the symbol name.
> > > > > > 
> > > > > > If name mangling can't be made identical (and it looks like it may be the case), we can probably work around it. I.e. for symbols that must have identical names on both sides we can generate a unique alias that's identical on both sides and use it instead when CUDA needs it.
> > > > > > 
> > > > > But why is LLVM responsible for mangling? Shouldn't this be done by the Clang frontend?
> > > > > I've found this test in `libcxxabi/test/test_demangle.pass.cpp`:
> > > > > ```
> > > > > {"_ZNK10__cxxabiv111__libcxxabi5__sub20first_demangled_nameEPc.eh", "__cxxabiv1::__libcxxabi::__sub::first_demangled_name(char*) const (.eh)"},
> > > > > ```
> > > > > 
> > > > > As said, I can't imagine a user defined value where LLVM needs to generate a unique name, so this should be fine for the CUDA functions you mentioned because they operate on variables.
> > > > Sorry. I should've said "the way LLVM creates unique names".  C++ mangling is indeed handled by clang.
> > > > 
> > > > AFAICT, we do need to consider possibility of LLVM generating unique name whenever  Clang calls LLVM's Value::setName(). The API explicitly guarantees that the name will be unique.
> > > > 
> > > > Another case would be an identifier with a unicode symbol in it. NVPTX would have to sanitize it with the result potentially clashing with existing name -- similar to @.str test case below. This is hypothetical at the moment as clang does not support unicode (yet?).
> > > > 
> > > > 
> > > All right, then the next natural question is: Do we need to keep the C++ mangling in mind when creating the unique name? Because that is currently the reason that there is a dot separator. If we don't need it we could just drop the dot and every target would be happy. In addition this would guarantee that LLVM will produce a "valid" name after NVPTX sanitized the name and there was a crash.
> > Alas dropping the dot does mess with the mangling. In that respect '$' would be lesser evil, at least we'll get the pre-uniquification C++ name demangled correctly. I'm not sure what effect that would have on demangling on windows.
> > ```
> > $ c++filt _Z1fv _Z1fv.1 _Z1fv$2 _Z1fv3
> > f()
> > f() [clone .1]
> > f()
> > _Z1fv3
> > 
> > $ bin/llvm-cxxfilt _Z1fv _Z1fv.1 _Z1fv$2 _Z1fv3
> > f()
> > f() (.1)
> > f()
> > _Z1fv3
> > ```
> > 
> > It looks like we'll have to use '$' for uniquification in nvptx as your patch does. That, at least, will deal with the part of the issue we need right now -- generate PTX which ptxas can compile and keep demanglers working. The cases where uniquification mismatch  would cause problems should be rare. We can solve that problem separately. It's not perfect, but it looks like we don't have many options here. 
> > 
> > I don't have any better ideas. 
> > @hfinkel, @rnk -- your thoughts? 
> Hehe, I guess you executed this on a shell? Then `$2` is an (empty) variable and the right commands are:
> ```
> $ c++filt _Z1fv _Z1fv.1 _Z1fv\$2 _Z1fv3
> f()
> f() [clone .1]
> _Z1fv$2
> _Z1fv3
> $ bin/llvm-cxxfilt _Z1fv _Z1fv.1 _Z1fv\$2 _Z1fv3
> f()
> f() (.1)
> _Z1fv$2
> _Z1fv3
> ```
> (at least the two demanglers agree...)
> 
> So a dollar sign doesn't help much which means we could also drop it entirely.
Oops. :-(
Well,  bummer. I guess we'll have to live with un-demangle-able unique names on NVPTX side. It's still an improvement over not being able to compile anything with such names.


https://reviews.llvm.org/D40573





More information about the llvm-commits mailing list