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

[Artem Belevich via Phabricator via llvm-commits]

tra added a comment.

There must be some truth in the saying "naming is one of the hardest problems in computer science". :-/



================
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 << ".";
+    }
----------------
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.




https://reviews.llvm.org/D40573