[SYCL] Avoid infinite loop when kernel fails to compile with memory error (#18888)

The runtime tries to build a kernel again if compilation fails. But if
UR returns a memory error the attempt counter was not compared against
the maximum number of attempts, so the compiler was continuously called
and eventually the loop counter would have overflowed.

Author: VerenaBeckham

sycl/source/detail/kernel_program_cache.hpp

@@ -818,7 +818,9 @@ class KernelProgramCache {
             BuildResult->Error.Code == UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY) {
           reset();
           BuildResult->updateAndNotify(BuildState::BS_Initial);
-          continue;
+          if (AttemptCounter + 1 < MaxAttempts) {
+            continue;
+          }
         }
 
         BuildResult->updateAndNotify(BuildState::BS_Failed);

sycl/unittests/kernel-and-program/OutOfResources.cpp

@@ -109,6 +109,31 @@ TEST_P(OutOfResourcesTestSuite, urProgramCreate) {
   }
 }
 
+static ur_result_t ProgramCreateWithILAlwaysFail(void *) { return ErrorCode; }
+
+TEST_P(OutOfResourcesTestSuite, urProgramCreateAlwaysFail) {
+  sycl::unittest::UrMock<> Mock;
+  ErrorCode = GetParam();
+  mock::getCallbacks().set_before_callback("urProgramCreateWithIL",
+                                           &ProgramCreateWithILAlwaysFail);
+
+  sycl::platform Plt{sycl::platform()};
+  sycl::context Ctx{Plt};
+  auto CtxImpl = detail::getSyclObjImpl(Ctx);
+  queue Q(Ctx, default_selector_v);
+
+  bool ThrewException = false;
+
+  try {
+    Q.single_task<class OutOfResourcesKernel1>([] {});
+  } catch (exception &Ex) {
+    auto Code = detail::get_ur_error(Ex);
+    EXPECT_EQ(Code, ErrorCode);
+    ThrewException = true;
+  }
+  EXPECT_TRUE(ThrewException);
+}
+
 static int nProgramLink = 0;
 
 static ur_result_t redefinedProgramLink(void *) {