only introduce __ldg() calls on accessed to global memory

__ldg() only makes sense on accesses to global memory.
It is not immediately obvious whether this would ever be introduced
on accesses to promoted tensor, i.e., whether accesses to read-only
tensors would ever get promoted, but it is better to be safe
than sorry and therefore to make sure __ldg() calls are never
introduced on accesses to promoted tensors.

Author: Sven Verdoolaege

tc/core/polyhedral/cuda/codegen.cc

@@ -398,9 +398,16 @@ struct LdgWrapper {
 
 template <typename AFF>
 void emitAccess(AFF access, const CodegenStatementContext& context) {
-  LdgWrapper ldgWrapper(context, access.get_tuple_id(isl::dim_type::out));
   context.ss << context.build().access_from(access).to_C_str();
 }
+
+// Print an access to global memory, wrapping the access in an "__ldg()"
+// call if the accessed tensor is known to be read-only.
+template <typename AFF>
+void emitGlobalAccess(AFF access, const CodegenStatementContext& context) {
+  LdgWrapper ldgWrapper(context, access.get_tuple_id(isl::dim_type::out));
+  emitAccess(access, context);
+}
 } // namespace
 
 void emitCopyStmt(const CodegenStatementContext& context) {
@@ -414,9 +421,9 @@ void emitCopyStmt(const CodegenStatementContext& context) {
   if (isRead) {
     emitAccess(isl::multi_pw_aff(promoted), context);
     context.ss << " = ";
-    emitAccess(isl::multi_pw_aff(original), context);
+    emitGlobalAccess(isl::multi_pw_aff(original), context);
   } else {
-    emitAccess(isl::multi_pw_aff(original), context);
+    emitGlobalAccess(isl::multi_pw_aff(original), context);
     context.ss << " = ";
     emitAccess(isl::multi_pw_aff(promoted), context);
   }