register promotion: insert syncs if scoped above thread mapping

Promotion to registers is performed in a particular scope.  Within this
scope, if a promotion is deemed valid, tensor elements are only accessed
by the same thread making thread synchronization unnecessary.  However,
outside this scope, promoted elements may be accessed by different
threads which would require synchronization. If different threads (or,
in practice, different iterations of thread-mapped loops) access the
same element in a way that requires synchronization, this is reflected
in the dependence relation.  The OuterBlockInnerThread mapping strategy
detects it and introduces synchronization statements above thread
mapping.  It only performs promotion to registers below thread mapping.
Therefore, synchronizations were unnecessary around copies to or from
registers.

PR #489 introduced functionality to promote to registers at any scope,
including above thread mapping.  In this case, synchronizations inserted
below may not suffice.  For example, in the tree of the shape

band(                  // contains a sequential loop
  // <- promotion scope
  extension(
    sequence(
      filter(          // main computation
        mapping(...)), // to threads
      filter(...)      // synchronization

different iterations of the outer sequential loop may lead to different
threads accessing the same tensor element (but in one iteration, only
one thread accesses it).  Copies from global memory to registers will be
inserted at the promotion scope, i.e. after the synchronization
statement.  A write to global memory by one thread will not be
synchronized with a read from the same address by a potentially
different thread in the following iteration of the loop above the
scoping point.  A synchronization must be inserted either before the
read from global memory or after the write-back.  In this particular
case, one may want to insert the write-back before the existing
synchronization, but it is not always possible in the general case where
the scoping point may be not immediately above the thread mapping.
Furthermore, it may also be necessary to synchronize due to dependences
with sibling subtrees that have a different mapping.

When register promotion copies are inserted above thread mapping,
introduce thread synchronizations before the copy to register and after
the copy from register.  This is a conservative approximation.  Exact
analysis would require analyzing dependences between an instance of the
scope and the rest of the elements and is left for future work.

Author: ftynse

tc/core/polyhedral/cuda/memory_promotion_heuristic.cc

@@ -698,6 +698,26 @@ void promoteToRegistersBelow(MappedScop& mscop, detail::ScheduleTree* scope) {
           partialSched);
     }
   }
+
+  // Return immediately if nothing was promoted.
+  if (scope->numChildren() == 0 ||
+      !matchOne(extension(sequence(any())), scope->child({0}))) {
+    return;
+  }
+
+  // If promoting above thread mapping, insert synchronizations.
+  // It is possible that promoted array elements are accessed by different
+  // threads outside the current scope (either in different iterations of the
+  // scope loops, or in sibling subtrees).  For now, always insert
+  // synchronizations, similarly to copies to shared memory.
+  //
+  // TODO: The exact check for sync insertion requires the dependences between
+  // the elements in the scope and those before/after the scope and a check if
+  // the dependent instances belong to the same thread.
+  auto ancestors = scope->ancestors(root);
+  if (functional::Filter(isMappingTo<mapping::ThreadId>, ancestors).empty()) {
+    scop.insertSyncsAroundSeqChildren(scope->child({0, 0}));
+  }
 }
 
 // Promote at the positions of the thread specific markers.

test/test_cuda_mapper_memory_promotion.cc

@@ -493,10 +493,12 @@ def fun(float(N,K) A, float(K,M) B, float(N,M) C) -> (O) {
 
     expectNoABCPromotion(code);
 
+    auto preCopySyncPos = code.find("__syncthreads()", oDeclPos);
     auto o00Pos = code.find("_O_0[0][0]");
     auto o10Pos = code.find("_O_0[1][0]");
     auto o20Pos = code.find("_O_0[2][0]");
     auto o30Pos = code.find("_O_0[3][0]");
+    auto postCopySyncPos = code.find("__syncthreads()", o30Pos);
 
     EXPECT_TRUE(o00Pos != std::string::npos)
         << "expected constant subscripts in _O_0";
@@ -506,6 +508,25 @@ def fun(float(N,K) A, float(K,M) B, float(N,M) C) -> (O) {
         << "expected constant subscripts in _O_0";
     EXPECT_TRUE(o30Pos != std::string::npos)
         << "expected constant subscripts in _O_0";
+
+    EXPECT_TRUE(preCopySyncPos != std::string::npos)
+        << "expected synchronization to be inserted";
+    EXPECT_TRUE(postCopySyncPos != std::string::npos)
+        << "expected synchronization to be inserted";
+
+    EXPECT_TRUE(
+        preCopySyncPos < o00Pos && preCopySyncPos < o10Pos &&
+        preCopySyncPos < o20Pos && preCopySyncPos < o30Pos)
+        << "expected synchronization before copies to registers";
+
+    o00Pos = code.find("_O_0[0][0]", postCopySyncPos);
+    o10Pos = code.find("_O_0[1][0]", postCopySyncPos);
+    o20Pos = code.find("_O_0[2][0]", postCopySyncPos);
+    o30Pos = code.find("_O_0[3][0]", postCopySyncPos);
+    EXPECT_TRUE(
+        o00Pos == std::string::npos && o10Pos == std::string::npos &&
+        o20Pos == std::string::npos && o20Pos == std::string::npos)
+        << "expected synchronization after copies from registers";
   }
 };