[SYCL] Add exception on recompile of AOT compiled kernel by program api

Signed-off-by: Alexander Flegontov <[email protected]>

Author: alexanderfle

sycl/source/detail/program_impl.cpp

@@ -231,7 +231,7 @@ void program_impl::build_with_kernel_name(string_class KernelName,
     if (is_cacheable_with_options(BuildOptions)) {
       MProgramAndKernelCachingAllowed = true;
       MProgram = ProgramManager::getInstance().getBuiltPIProgram(
-          Module, get_context(), KernelName, this);
+          Module, get_context(), KernelName, this, true);
       const detail::plugin &Plugin = getPlugin();
       Plugin.call<PiApiKind::piProgramRetain>(MProgram);
     } else {
@@ -440,7 +440,7 @@ void program_impl::create_pi_program_with_kernel_name(
   assert(!MProgram && "This program already has an encapsulated PI program");
   ProgramManager &PM = ProgramManager::getInstance();
   RTDeviceBinaryImage &Img =
-      PM.getDeviceImage(Module, KernelName, get_context());
+      PM.getDeviceImage(Module, KernelName, get_context(), true);
   MProgram = PM.createPIProgram(Img, get_context());
 }
 

sycl/source/detail/program_manager/program_manager.cpp

@@ -136,13 +136,13 @@ static RT::PiProgram createSpirvProgram(const ContextImplPtr Context,
 
 RTDeviceBinaryImage &
 ProgramManager::getDeviceImage(OSModuleHandle M, const string_class &KernelName,
-                               const context &Context) {
+                               const context &Context, bool FromProgramApi) {
   if (DbgProgMgr > 0)
     std::cerr << ">>> ProgramManager::getDeviceImage(" << M << ", \""
               << KernelName << "\", " << getRawSyclObjImpl(Context) << ")\n";
 
   KernelSetId KSId = getKernelSetId(M, KernelName);
-  return getDeviceImage(M, KSId, Context);
+  return getDeviceImage(M, KSId, Context, FromProgramApi);
 }
 
 template <typename ExceptionT, typename RetT>
@@ -360,7 +360,7 @@ RT::PiProgram ProgramManager::createPIProgram(const RTDeviceBinaryImage &Img,
 RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
                                                 const context &Context,
                                                 const string_class &KernelName,
-                                                const program_impl *Prg) {
+                                                const program_impl *Prg, bool FromProgramApi) {
   KernelSetId KSId = getKernelSetId(M, KernelName);
 
   const ContextImplPtr Ctx = getSyclObjImpl(Context);
@@ -376,8 +376,8 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
   auto GetF = [](const Locked<ProgramCacheT> &LockedCache) -> ProgramCacheT& {
     return LockedCache.get();
   };
-  auto BuildF = [this, &M, &KSId, &Context, Prg] {
-    const RTDeviceBinaryImage &Img = getDeviceImage(M, KSId, Context);
+  auto BuildF = [this, &M, &KSId, &Context, Prg, &FromProgramApi] {
+    const RTDeviceBinaryImage &Img = getDeviceImage(M, KSId, Context, FromProgramApi);
 
     ContextImplPtr ContextImpl = getSyclObjImpl(Context);
     const detail::plugin &Plugin = ContextImpl->getPlugin();
@@ -634,7 +634,8 @@ ProgramManager::ProgramManager() {
 
 RTDeviceBinaryImage &ProgramManager::getDeviceImage(OSModuleHandle M,
                                                     KernelSetId KSId,
-                                                    const context &Context) {
+                                                    const context &Context,
+                                                    bool FromProgramApi) {
   if (DbgProgMgr > 0) {
     std::cerr << ">>> ProgramManager::getDeviceImage(" << M << ", \"" << KSId
               << "\", " << getRawSyclObjImpl(Context) << ")\n";
@@ -660,6 +661,18 @@ RTDeviceBinaryImage &ProgramManager::getDeviceImage(OSModuleHandle M,
 
   Ctx->getPlugin().call<PiApiKind::piextDeviceSelectBinary>(
       getFirstDevice(Ctx), RawImgs.data(), (cl_uint)RawImgs.size(), &ImgInd);
+
+  if (FromProgramApi) {
+    // if the image is already compiled with AOT then throws exception
+    const pi_device_binary_struct &RawImg = Imgs[ImgInd]->getRawData();
+    if ((strcmp(RawImg.DeviceTargetSpec, PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64) == 0) ||
+        (strcmp(RawImg.DeviceTargetSpec, PI_DEVICE_BINARY_TARGET_SPIRV64_GEN) == 0) ||
+        (strcmp(RawImg.DeviceTargetSpec, PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA) == 0)) {
+          throw feature_not_supported("Recompiling AOT image is not supported",
+                                      PI_INVALID_OPERATION);
+    }
+  }
+
   Img = Imgs[ImgInd].get();
 
   if (DbgProgMgr > 0) {

sycl/source/detail/program_manager/program_manager.hpp

@@ -61,7 +61,8 @@ class ProgramManager {
   static ProgramManager &getInstance();
   RTDeviceBinaryImage &getDeviceImage(OSModuleHandle M,
                                       const string_class &KernelName,
-                                      const context &Context);
+                                      const context &Context,
+                                      bool FromProgramApi = false);
   RT::PiProgram createPIProgram(const RTDeviceBinaryImage &Img,
                                 const context &Context);
   /// Builds or retrieves from cache a program defining the kernel with given
@@ -76,7 +77,7 @@ class ProgramManager {
   ///        once the function returns.
   RT::PiProgram getBuiltPIProgram(OSModuleHandle M, const context &Context,
                                   const string_class &KernelName,
-                                  const program_impl *Prg = nullptr);
+                                  const program_impl *Prg = nullptr, bool FromProgramApi = false);
   RT::PiKernel getOrCreateKernel(OSModuleHandle M, const context &Context,
                                  const string_class &KernelName,
                                  const program_impl *Prg);
@@ -110,7 +111,7 @@ class ProgramManager {
   ProgramManager &operator=(ProgramManager const &) = delete;
 
   RTDeviceBinaryImage &getDeviceImage(OSModuleHandle M, KernelSetId KSId,
-                                      const context &Context);
+                                      const context &Context, bool FromProgramApi = false);
   using ProgramPtr = unique_ptr_class<remove_pointer_t<RT::PiProgram>,
                                       decltype(&::piProgramRelease)>;
   ProgramPtr build(ProgramPtr Program, const ContextImplPtr Context,

sycl/test/aot/spec_const_aot.cpp

@@ -38,7 +38,6 @@ int main(int argc, char **argv) {
   cl::sycl::experimental::spec_constant<int32_t, MyInt32Const> i32 =
       prog.set_spec_constant<MyInt32Const>(10);
 
-  prog.build_with_kernel_type<Kernel>();
 
   std::vector<int> vec(1);
   {
@@ -47,7 +46,6 @@ int main(int argc, char **argv) {
     q.submit([&](cl::sycl::handler &cgh) {
       auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
       cgh.single_task<Kernel>(
-          prog.get_kernel<Kernel>(),
           [=]() {
             acc[0] = i32.get();
           });