Refactor may-have-IGBAs -> has-no-IGBAs

Author: Henry Linjamäki

llvm_passes/HipIGBADetector.cpp

@@ -10,21 +10,23 @@
 // buffer accesses (IGBA). The outcome of the analysis is stored in a magic
 // variable for the chipStar runtime:
 //
-//   uint8_t __chip_module_may_have_IGBAs = <result>;
+//   uint8_t __chip_module_has_no_IGBAs = <result>;
 //
-// Where the result is one if the are potential IGBAs and otherwise zero.
+// Where the result is one if the are no potential IGBAs and otherwise it is
+// zero.
 //
-// If there would be an IGPA in the module, there has to to be a load instruction
-// with a pointer operand which is either loaded from memory or crafted from an
-// integer (which OTOH is loaded from somewhere else). The analysis is very
-// simple and naive: we look for pointer load and inttoptr instructions in the
-// whole module. If we see any, we conclude there are potential IGBAs. Downsides
-// of this are that
+// If there would be an IGPA in the module, there has to to be a load
+// instruction with a pointer operand which is either loaded from memory or
+// crafted from an integer (which OTOH is loaded from somewhere else). The
+// analysis is very simple and naive: we look for pointer load and inttoptr
+// instructions in the whole module. If we see any, we conclude there are
+// potential IGBAs. Downsides of this are that
 //
-// * may-have-IGBAs flag is raised even tough only one kernel has IGBAs an
+// * may-have-IGBAs is concluded even tough only one kernel has IGBAs an
 //   others don't
 //
-// * unoptimized modules (-O0) will likely raise may-have-IGBAs flag.
+// * unoptimized modules (-O0) will likely likely result in may-have-IGBAs
+// * conclusion
 //
 // The motivation for this analysis is to reduce clSetKernelExecInfo() calls in
 // the OpenCL backend.
@@ -58,15 +60,15 @@ static bool hasPotentialIGBAs(Module &M) {
 }
 
 static bool detectIGBAs(Module &M) {
-  constexpr auto *MagicVarName = "__chip_module_may_have_IGBAs";
+  constexpr auto *MagicVarName = "__chip_module_has_no_IGBAs";
 
   if (M.getGlobalVariable(MagicVarName))
     return false; // Bail out: the module has already been processed.
 
   bool Result = hasPotentialIGBAs(M);
   LLVM_DEBUG(dbgs() << "Has IGBAs: " << Result << "\n");
 
-  auto *Init = ConstantInt::get(IntegerType::get(M.getContext(), 8), Result);
+  auto *Init = ConstantInt::get(IntegerType::get(M.getContext(), 8), !Result);
   (void)new GlobalVariable(
       M, Init->getType(), true,
       // Mark the GV as external for keeping it alive at least until the

src/backend/OpenCL/CHIPBackendOpenCL.cc

@@ -235,10 +235,9 @@ annotateIndirectPointers(const CHIPContextOpenCL &Ctx,
   // If we have determined that the module does not have indirect
   // global memory accesses (IGBAs; see HipIGBADetectorPass), we may
   // skip the annotation.
-  if (!ModInfo.MayHaveIGBAs)
+  if (ModInfo.HasNoIGBAs)
     return nullptr;
 
-
   std::unique_ptr<std::vector<std::shared_ptr<void>>> AllocKeepAlives;
   std::vector<void *> AnnotationList;
   LOCK(Ctx.ContextMtx); // CHIPContextOpenCL::MemManager_

src/common.hh

@@ -57,9 +57,9 @@ using SPVFunctionInfoMap = std::map<std::string, std::shared_ptr<SPVFuncInfo>>;
 struct SPVModuleInfo {
   SPVFunctionInfoMap FuncInfoMap;
 
-  /// Set to true if the module may have indirect global buffer
-  /// accesses (IGBA) in any kernel.
-  bool MayHaveIGBAs = true;
+  /// Set to true if the module is known not to have indirect global
+  /// buffer accesses (IGBA) in any kernel.
+  bool HasNoIGBAs = false;
 };
 
 bool filterSPIRV(const char *Bytes, size_t NumBytes, std::string &Dst);

src/spirv.cc

@@ -516,11 +516,11 @@ class SPIRVmodule {
   std::map<std::string_view, std::vector<std::pair<uint16_t, uint16_t>>>
       SpilledArgAnnotations_;
 
-  // This flag indicates if the module may have indirect global buffer
-  // accesses (IGBA) in any kernel. This is told by a magic variable
-  // created by HipIGBADetectorPass. Defaults to true in case the
-  // variable is not found.
-  bool MayHaveIGBAs_ = true;
+  // This flag indicates if the module is known not to have indirect
+  // global buffer accesses (IGBA) in any kernel. This is told by a
+  // magic variable created by HipIGBADetectorPass. Defaults to false
+  // in case the variable is not found.
+  bool HasNoIGBAs_ = false;
 
   bool MemModelCL_;
   bool KernelCapab_;
@@ -587,7 +587,7 @@ class SPIRVmodule {
     }
     KernelInfoMap_.clear();
 
-    ModuleInfo.MayHaveIGBAs = MayHaveIGBAs_;
+    ModuleInfo.HasNoIGBAs = HasNoIGBAs_;
 
     return true;
   }
@@ -733,11 +733,11 @@ class SPIRVmodule {
         }
 
         // A magic variable created by HipIGBADetector.cpp.
-        if (Name == "__chip_module_may_have_IGBAs") {
+        if (Name == "__chip_module_has_no_IGBAs") {
           // Get initializer operand.
           auto *Init = getInstruction(Inst->getWord(4));
           // Init is known to be 8-bit unsigned constant.
-          MayHaveIGBAs_ = Init->getWord(3);
+          HasNoIGBAs_ = Init->getWord(3);
         }
       }
 

tests/runtime/CMakeLists.txt

@@ -102,6 +102,6 @@ target_compile_definitions(TestBitInsert
 
 add_hip_runtime_test(TestBallot.hip)
 
-add_hip_runtime_test(TestPositiveIGBADetection1.hip)
-add_hip_runtime_test(TestPositiveIGBADetection2.hip)
-add_hip_runtime_test(TestNegativeIGBADetection.hip)
+add_hip_runtime_test(TestNegativeHasNoIGBAs1.hip)
+add_hip_runtime_test(TestNegativeHasNoIGBAs2.hip)
+add_hip_runtime_test(TestPositiveHasNoIGBAs.hip)

…s/runtime/TestPositiveIGBADetection1.hip tests/runtime/TestNegativeHasNoIGBAs1.hiptests/runtime/TestPositiveIGBADetection1.hip renamed to tests/runtime/TestNegativeHasNoIGBAs1.hip tests/runtime/TestPositiveIGBADetection1.hip renamed to tests/runtime/TestNegativeHasNoIGBAs1.hip

@@ -19,7 +19,6 @@ __global__ void k(SomeStruct Args) { *Args.Ptr = 123; }
 int main() {
   auto *ChipKernel = Backend->getActiveDevice()->findKernel(
       HostPtr(reinterpret_cast<const void *>(k)));
-  if (ChipKernel->getModule()->getInfo().MayHaveIGBAs)
-    return 0;
-  return 1;
+  // HasNoIGBAs should be false here.
+  return ChipKernel->getModule()->getInfo().HasNoIGBAs;
 }

…s/runtime/TestPositiveIGBADetection2.hip tests/runtime/TestNegativeHasNoIGBAs2.hiptests/runtime/TestPositiveIGBADetection2.hip renamed to tests/runtime/TestNegativeHasNoIGBAs2.hip tests/runtime/TestPositiveIGBADetection2.hip renamed to tests/runtime/TestNegativeHasNoIGBAs2.hip

@@ -17,7 +17,6 @@ __global__ void k(uintptr_t Arg) {
 int main() {
   auto *ChipKernel = Backend->getActiveDevice()->findKernel(
       HostPtr(reinterpret_cast<const void *>(k)));
-  if (ChipKernel->getModule()->getInfo().MayHaveIGBAs)
-    return 0;
-  return 1;
+  // HasNoIGBAs should be false here.
+  return ChipKernel->getModule()->getInfo().HasNoIGBAs;
 }

…ts/runtime/TestNegativeIGBADetection.hip tests/runtime/TestPositiveHasNoIGBAs.hiptests/runtime/TestNegativeIGBADetection.hip renamed to tests/runtime/TestPositiveHasNoIGBAs.hip tests/runtime/TestNegativeIGBADetection.hip renamed to tests/runtime/TestPositiveHasNoIGBAs.hip

@@ -15,7 +15,6 @@ __global__ void k(int *Out) { *Out = 123; }
 int main() {
   auto *ChipKernel = Backend->getActiveDevice()->findKernel(
       HostPtr(reinterpret_cast<const void *>(k)));
-  if (!ChipKernel->getModule()->getInfo().MayHaveIGBAs)
-    return 0;
-  return 1;
+  // HasNoIGBAs should be true here.
+  return !ChipKernel->getModule()->getInfo().HasNoIGBAs;
 }