gpu: jit: disable mid-thread preemption

petercad · petercad · commit 89d6f876479c · 2020-10-09T13:40:14.000-07:00
diff --git a/src/gpu/jit/ngen/ngen_interface.hpp b/src/gpu/jit/ngen/ngen_interface.hpp
@@ -257,6 +257,8 @@ void NEOInterfaceHandler::generateDummyCL(std::ostream &stream) const
     }
     stream << ") {\n";
     stream << "    global volatile int *____;\n";
+    if (hw == HW::Gen9)
+        stream << "    volatile double *__df; *__df = 1.1 / *__df;\n";
 
     if (needLocalID)        stream << "    (void) ____[get_local_id(0)];\n";
     if (needLocalSize)      stream << "    (void) ____[get_enqueued_local_size(0)];\n";

Original file line number	Diff line number	Diff line change
`@@ -257,6 +257,8 @@ void NEOInterfaceHandler::generateDummyCL(std::ostream &stream) const`
`257`	`257`	`}`
`258`	`258`	`stream << ") {\n";`
`259`	`259`	`stream << " global volatile int *____;\n";`
	`260`	`+ if (hw == HW::Gen9)`
	`261`	`+ stream << " volatile double __df; __df = 1.1 / *__df;\n";`
`260`	`262`
`261`	`263`	`if (needLocalID) stream << " (void) ____[get_local_id(0)];\n";`
`262`	`264`	`if (needLocalSize) stream << " (void) ____[get_enqueued_local_size(0)];\n";`