gpu: jit: ngen: prepare for proposed IGC changes

oneapi-src · Nov 2, 2022 · 7144393 · 7144393
1 parent d23cc95
commit 7144393
Showing 1 changed file with 6 additions and 0 deletions.
diff --git a/src/gpu/jit/ngen/ngen_interface.hpp b/src/gpu/jit/ngen/ngen_interface.hpp
@@ -349,6 +349,12 @@ void InterfaceHandler::generateDummyCL(std::ostream &stream) const
             stream << "    __asm__ volatile(\"\" :: \"rw.u\"(" << assignment.name << "));\n";
     }
 
+    for (const auto &assignment : assignments) {
+        // Force IGC to assume surface accesses could occur if necessary.
+        if (assignment.exttype == ExternalArgumentType::GlobalPtr && assignment.globalSurfaceAccess())
+            stream << "    { volatile uchar __load = ((global uchar *) " << assignment.name << ")[get_local_id(0)];}\n";
+    }
+
     stream << "}\n";
 }