add wrapper around __ldg for those poor souls with a very old CUDA card

__ldg is apparently only supported since 3.5.  Add a wrapper to ignore
the __ldg call on older architectures.
The wrapper is put inside the __tc namespace in analogy
with __tc::CubReduceAlongX, even though names that start
with a double underscore are supposed to be reserved.

The wrapper was inspired by https://stackoverflow.com/a/27302007

Author: Sven Verdoolaege

tc/core/libraries.h

@@ -156,6 +156,25 @@ template<typename T> inline __device__ T floord(T n, T d) {
 } // namespace cpp
 
 namespace cuda {
+
+// Wrapper around __ldg to avoid compilation errors
+// on architectures that do not support it.
+constexpr auto ldg = R"CUDA(
+
+namespace __tc {
+template<typename T>
+__device__ __forceinline__ T ldg(const T* ptr) {
+#if __CUDA_ARCH__ >= 350
+  return __ldg(ptr);
+#else
+  return *ptr;
+#endif
+}
+} // namespace __tc
+)CUDA";
+
+const static std::string kLdg = "__tc::ldg";
+
 constexpr auto common = R"CUDA(
 
 namespace __tc {

tc/core/polyhedral/cuda/codegen.cc

@@ -381,7 +381,7 @@ struct LdgWrapper {
   LdgWrapper(const CodegenStatementContext& context, isl::id id)
       : readOnly_(context.readOnlySet.count(id) > 0), out_(context.ss) {
     if (readOnly_) {
-      out_ << "__ldg(&";
+      out_ << tc::code::cuda::kLdg << "(&";
     }
   }
 

tc/core/polyhedral/cuda/mapped_scop.cc

@@ -919,6 +919,9 @@ std::tuple<std::string, tc::Grid, tc::Block> MappedScop::codegen(
   code << code::cpp::boundsAsTemplate << code::c::types << code::c::defines;
   code << code::c::warpSyncFunctions;
   code << std::endl;
+  if (useReadOnlyCache) {
+    code << code::cuda::ldg;
+  }
   if (mappedScopForCodegen->scop().treeSyncUpdateMap.size() != 0) {
     code << code::cuda::common;
     code << code::cuda::cubBlockReduce;

test/test_cuda_mapper.cc

@@ -1018,8 +1018,9 @@ def fun(float(N) I) -> (O) {
 )TC";
   auto mappingOptions = DefaultOptions().useReadOnlyCache(true);
   auto code = codegenMapped(tc, mappingOptions);
-  ASSERT_TRUE(code.find("__ldg(&O") == std::string::npos) << code; // no
-  ASSERT_TRUE(code.find("__ldg(&I") != std::string::npos) << code; // yes
+  using tc::code::cuda::kLdg;
+  ASSERT_TRUE(code.find(kLdg + "(&O") == std::string::npos) << code; // no
+  ASSERT_TRUE(code.find(kLdg + "(&I") != std::string::npos) << code; // yes
 }
 
 /*