[SYCL][E2E] fixed UR FileCheck conditions to match new traces

makes llvm work with this change: oneapi-src/unified-runtime#2101
fixes this bug: oneapi-src/unified-runtime#2002

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -116,14 +116,8 @@ if(SYCL_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 7384e2d7b908de0b2bca9f3c57827ea84698864e
-  # Merge: ed85c3e2 a32549b4
-  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
-  # Date:   Fri Sep 20 14:56:40 2024 +0100
-  #   Merge pull request #2108 from pbalcer/fix-perf-tracing-logs
-  #   fix perf regression in tracing layer
-  set(UNIFIED_RUNTIME_TAG 7384e2d7b908de0b2bca9f3c57827ea84698864e)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/lslusarczyk/unified-runtime.git")
+  set(UNIFIED_RUNTIME_TAG 12e0cee8f7ef2761c1ce8d45312855dc84140ee0)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

sycl/test-e2e/Basic/buffer/native_buffer_creation_flags.cpp

@@ -20,7 +20,7 @@ int main() {
   Q.submit([&](handler &Cgh) {
     // Now that we have a read-write host allocation, check that the native
     // buffer is created with the UR_MEM_FLAG_USE_HOST_POINTER flag.
-    // CHECK: urMemBufferCreate
+    // CHECK: ---> urMemBufferCreate
     // CHECK-SAME: UR_MEM_FLAG_USE_HOST_POINTER
     auto BufAcc = Buf.get_access<access::mode::read>(Cgh);
     Cgh.single_task<Foo>([=]() { int A = BufAcc[0]; });

sycl/test-e2e/Basic/buffer/subbuffer_overlap.cpp

@@ -16,13 +16,13 @@ int main() {
   for (auto &e : sycl::host_accessor{b})
     e = idx++ % size;
 
-  // CHECK: urMemBufferPartition
+  // CHECK: ---> urMemBufferPartition
   // CHECK: .origin = 256, .size = 64
   q.submit([&](sycl::handler &cgh) {
     sycl::accessor acc{sub1, cgh};
     cgh.parallel_for(size, [=](auto id) { acc[id] += 1; });
   });
-  // CHECK: urMemBufferPartition
+  // CHECK: ---> urMemBufferPartition
   // CHECK: .origin = 256, .size = 128
   q.submit([&](sycl::handler &cgh) {
     sycl::accessor acc{sub2, cgh};

sycl/test-e2e/Basic/event_release.cpp

@@ -29,8 +29,8 @@ int main() {
   // Buffer destruction triggers execution graph cleanup, check that both
   // events (one for launching the kernel and one for memory transfer to host)
   // are released.
-  // CHECK: urEventRelease
-  // CHECK: urEventRelease
+  // CHECK: ---> urEventRelease
+  // CHECK: ---> urEventRelease
   assert(Val == Gold);
   // CHECK: End of main scope
   std::cout << "End of main scope" << std::endl;

sycl/test-e2e/Basic/fill_accessor_ur.cpp

@@ -129,19 +129,19 @@ int main() {
 }
 
 // CHECK: start testFill_Buffer1D
-// CHECK: urEnqueueMemBufferFill
+// CHECK: ---> urEnqueueMemBufferFill
 // CHECK: start testFill_Buffer1D -- OFFSET
-// CHECK: urEnqueueMemBufferFill
+// CHECK: ---> urEnqueueMemBufferFill
 
 // CHECK: start testFill_Buffer2D
-// CHECK: urEnqueueMemBufferFill
+// CHECK: ---> urEnqueueMemBufferFill
 // CHECK: start testFill_Buffer2D -- OFFSET
-// CHECK: urEnqueueKernelLaunch
+// CHECK: ---> urEnqueueKernelLaunch
 
 // CHECK: start testFill_Buffer3D
-// CHECK: urEnqueueMemBufferFill
+// CHECK: ---> urEnqueueMemBufferFill
 // CHECK: start testFill_Buffer3D -- OFFSET
-// CHECK: urEnqueueKernelLaunch
+// CHECK: ---> urEnqueueKernelLaunch
 
 // CHECK: start testFill_ZeroDim
-// CHECK: urEnqueueMemBufferFill
+// CHECK: ---> urEnqueueMemBufferFill

sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp

@@ -150,26 +150,26 @@ int main() {
 
     sycl::kernel_bundle<sycl::bundle_state::object> KernelBundleObject1 =
         sycl::compile(KernelBundleInput1, KernelBundleInput1.get_devices());
-    // CHECK:---> urProgramCreate
+    // CHECK:<--- urProgramCreate
     // CHECK-SAME:, .phProgram = {{.*}} ([[PROGRAM_HANDLE1:[0-9a-fA-Fx]+]])
     // CHECK-SAME: -> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urProgramCompile
+    // CHECK:<--- urProgramCompile
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]]
 
     sycl::kernel_bundle<sycl::bundle_state::object> KernelBundleObject2 =
         sycl::compile(KernelBundleInput2, KernelBundleInput2.get_devices());
-    // CHECK:---> urProgramCreate
+    // CHECK:<--- urProgramCreate
     // CHECK-SAME:, .phProgram = {{.*}} ([[PROGRAM_HANDLE2:[0-9a-fA-Fx]+]])
     // CHECK-SAME: -> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urProgramCompile(
+    // CHECK:<--- urProgramCompile(
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE2]]
 
     sycl::kernel_bundle<sycl::bundle_state::executable> KernelBundleExecutable =
         sycl::link({KernelBundleObject1, KernelBundleObject2},
                    KernelBundleObject1.get_devices());
-    // CHECK:---> urProgramLink{{.*}} -> UR_RESULT_SUCCESS;
+    // CHECK:<--- urProgramLink{{.*}} -> UR_RESULT_SUCCESS;
     // UR tracing doesn't allow checking for all input programs so far.
 
     assert(KernelBundleExecutable.has_kernel(Kernel1ID));
@@ -179,14 +179,14 @@ int main() {
         KernelBundleExecutable2 =
             sycl::build(KernelBundleInput1, KernelBundleInput1.get_devices());
 
-    // CHECK:---> urProgramCreate
+    // CHECK:<--- urProgramCreate
     // CHECK-SAME:, .phProgram = {{.*}} ([[PROGRAM_HANDLE3:[0-9a-fA-Fx]+]])
     // CHECK-SAME: -> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urProgramBuild(
+    // CHECK:<--- urProgramBuild(
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE3]]
     //
-    // CHECK:---> urProgramRetain(
+    // CHECK:<--- urProgramRetain(
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE3]]
     // CHECK-SAME:-> UR_RESULT_SUCCESS;
 
@@ -204,31 +204,31 @@ int main() {
     sycl::kernel_bundle KernelBundleExecutable =
         sycl::get_kernel_bundle<sycl::bundle_state::executable>(Ctx, {Dev},
                                                                 {Kernel3ID});
-    // CHECK:---> urProgramCreate
+    // CHECK:<--- urProgramCreate
     // CHECK-SAME:, .phProgram = {{.*}} ([[PROGRAM_HANDLE4:[0-9a-fA-Fx]+]])
     // CHECK-SAME: -> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urProgramBuild(
+    // CHECK:<--- urProgramBuild(
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE4]]
     //
-    // CHECK:---> urProgramRetain(
+    // CHECK:<--- urProgramRetain(
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE4]]
     // CHECK-SAME:-> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urKernelCreate(
+    // CHECK:<--- urKernelCreate(
     // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE4]]
     // CHECK-SAME: .pKernelName = {{[0-9a-fA-Fx]+}} (_ZTS11Kernel3Name)
     // CHECK-SAME: .phKernel = {{[0-9a-fA-Fx]+}}  ([[KERNEL_HANDLE:[0-9a-fA-Fx]+]])
     // CHECK-SAME: -> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urKernelRetain(
+    // CHECK:<--- urKernelRetain(
     // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
     // CHECK-SAME:-> UR_RESULT_SUCCESS;
     //
-    // CHECK:---> urEnqueueKernelLaunch(
+    // CHECK:<--- urEnqueueKernelLaunch(
     // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
     //
-    // CHECK:---> urKernelRelease(
+    // CHECK:<--- urKernelRelease(
     // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
     // CHECK-SAME:-> UR_RESULT_SUCCESS;
 

sycl/test-e2e/DeviceCodeSplit/grf.cpp

@@ -145,10 +145,10 @@ int main(void) {
 
 // CHECK: ---> urKernelCreate({{.*}}SingleGRF{{.*}}-> UR_RESULT_SUCCESS
 
-// CHECK-NO-VAR: urProgramBuild{{.*}}-ze-opt-large-register-file
-// CHECK-WITH-VAR: urProgramBuild{{.*}}-g -ze-opt-large-register-file
-// CHECK-AUTO-NO-VAR: urProgramBuild{{.*}}-ze-intel-enable-auto-large-GRF-mode
-// CHECK-AUTO-WITH-VAR: urProgramBuild{{.*}}-g -ze-intel-enable-auto-large-GRF-mode
+// CHECK-NO-VAR: ---> urProgramBuild{{.*}}-ze-opt-large-register-file
+// CHECK-WITH-VAR: ---> urProgramBuild{{.*}}-g -ze-opt-large-register-file
+// CHECK-AUTO-NO-VAR: ---> urProgramBuild{{.*}}-ze-intel-enable-auto-large-GRF-mode
+// CHECK-AUTO-WITH-VAR: ---> urProgramBuild{{.*}}-g -ze-intel-enable-auto-large-GRF-mode
 // CHECK-SAME: -> UR_RESULT_SUCCESS
 
 // CHECK: ---> urKernelCreate({{.*}}SpecifiedGRF{{.*}}-> UR_RESULT_SUCCESS

sycl/test-e2e/DiscardEvents/discard_events_usm.cpp

@@ -71,7 +71,7 @@
 // CHECK: ---> urEnqueueUSMPrefetch
 // CHECK: .phEvent = nullptr
 //
-// CHECK: ---> urEnqueueUSMAdvise
+// CHECK: <--- urEnqueueUSMAdvise
 // CHECK: .phEvent = nullptr
 // CHECK-SAME: ) -> {{UR_RESULT_SUCCESS|UR_RESULT_ERROR_ADAPTER_SPECIFIC}}
 //

sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp

@@ -94,4 +94,4 @@ int main(void) {
 // Don't use -NEXT here to split the line because we need to allow for the
 // possbility of a BuildExp( that fails with UNSUPPORTED followed by a Build(
 // that results in SUCCESS
-// CHECK: ---> urProgramBuild{{(Exp)?}}({{.*}}-vc-codegen{{.*}} -> UR_RESULT_SUCCESS
+// CHECK: <--- urProgramBuild{{(Exp)?}}({{.*}}-vc-codegen{{.*}} -> UR_RESULT_SUCCESS

sycl/test-e2e/Graph/Explicit/kernel_bundle.cpp

@@ -5,27 +5,27 @@
 // Checks the UR call trace to ensure that the bundle kernel of the single task
 // is used.
 
-// CHECK:---> urProgramCreateWithIL(
+// CHECK:<--- urProgramCreateWithIL(
 // CHECK-SAME: .phProgram = {{.*}} ([[PROGRAM_HANDLE1:[0-9a-fA-Fx]+]])
 
 //
-// CHECK:---> urProgramBuildExp(
+// CHECK:<--- urProgramBuildExp(
 // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]]
 //
-// CHECK:---> urProgramRetain(.hProgram = [[PROGRAM_HANDLE1]]) -> UR_RESULT_SUCCESS
+// CHECK:<--- urProgramRetain(.hProgram = [[PROGRAM_HANDLE1]]) -> UR_RESULT_SUCCESS
 
-// CHECK:---> urKernelCreate(
+// CHECK:<--- urKernelCreate(
 // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]]
 // CHECK-SAME: .pKernelName = {{.*}} (_ZTS11Kernel1Name)
 // CHECK-SAME: .phKernel = {{.*}} ([[KERNEL_HANDLE:[0-9a-fA-Fx]+]])
 // CHECK-SAME: -> UR_RESULT_SUCCESS
 //
-// CHECK:---> urKernelRetain(.hKernel = [[KERNEL_HANDLE]]) -> UR_RESULT_SUCCESS
+// CHECK:<--- urKernelRetain(.hKernel = [[KERNEL_HANDLE]]) -> UR_RESULT_SUCCESS
 //
-// CHECK:---> urCommandBufferAppendKernelLaunchExp(
+// CHECK:<--- urCommandBufferAppendKernelLaunchExp(
 // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
 //
-// CHECK:---> urKernelRelease(.hKernel = [[KERNEL_HANDLE]]) -> UR_RESULT_SUCCESS
+// CHECK:<--- urKernelRelease(.hKernel = [[KERNEL_HANDLE]]) -> UR_RESULT_SUCCESS
 
 #define GRAPH_E2E_EXPLICIT
 

sycl/test-e2e/Graph/Explicit/memadvise.cpp

@@ -10,14 +10,14 @@
 // impact results but only performances, we verify
 // that a node is correctly added by checking UR function calls.
 
-// CHECK: urCommandBufferAppendUSMAdviseExp
+// CHECK: ---> urCommandBufferAppendUSMAdviseExp
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER:]]
 // CHECK-SAME: .pMemory = 0x[[#%x,PTR:]]
 // CHECK-SAME: .size = 400
 // CHECK-SAME: .pSyncPoint = {{.*}} (0x[[#%x,ADVISE_SYNC_POINT:]])
 // CHECK-SAME: -> UR_RESULT_SUCCESS
 
-// CHECK: urCommandBufferAppendKernelLaunchExp(
+// CHECK: ---> urCommandBufferAppendKernelLaunchExp(
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER]]
 // CHECK-SAME: .hKernel = 0x[[#%x,KERNEL:]]
 // CHECK-SAME: .workDim = 1

sycl/test-e2e/Graph/Explicit/prefetch.cpp

@@ -10,13 +10,13 @@
 // impact results but only performances, we verify
 // that a node is correctly added by checking UR function calls
 
-// CHECK: urCommandBufferAppendUSMPrefetchExp(
+// CHECK: ---> urCommandBufferAppendUSMPrefetchExp(
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER:]]
 // CHECK-SAME: .pMemory = 0x[[#%x,PTR:]]
 // CHECK-SAME: .size = 400
 // CHECK-SAME: .pSyncPoint = {{.*}} (0x[[#%x,PREFETCH_SYNC_POINT:]])
 
-// CHECK: urCommandBufferAppendKernelLaunchExp(
+// CHECK: ---> urCommandBufferAppendKernelLaunchExp(
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER]]
 // CHECK-SAME: .hKernel = 0x[[#%x,KERNEL:]]
 // CHECK-SAME: .workDim = 1

sycl/test-e2e/Graph/RecordReplay/kernel_bundle.cpp

@@ -5,31 +5,31 @@
 // Checks the UR call trace to ensure that the bundle kernel of the single task
 // is used.
 
-// CHECK:---> urProgramCreateWithIL(
+// CHECK:<--- urProgramCreateWithIL(
 // CHECK-SAME: .phProgram = {{.*}} ([[PROGRAM_HANDLE1:[0-9a-fA-Fx]+]])
 // CHECK-SAME: -> UR_RESULT_SUCCESS;
 //
-// CHECK:---> urProgramBuildExp(
+// CHECK:<--- urProgramBuildExp(
 // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]]
 //
-// CHECK:---> urProgramRetain(
+// CHECK:<--- urProgramRetain(
 // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]]
 // CHECK-SAME: -> UR_RESULT_SUCCESS;
 
-// CHECK:---> urKernelCreate(
+// CHECK:<--- urKernelCreate(
 // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]]
 // CHECK-SAME: .pKernelName = {{[0-9a-fA-Fx]+}} (_ZTS11Kernel1Name)
 // CHECK-SAME: .phKernel = {{[0-9a-fA-Fx]+}}  ([[KERNEL_HANDLE:[0-9a-fA-Fx]+]])
 // CHECK-SAME: -> UR_RESULT_SUCCESS;
 //
-// CHECK:---> urKernelRetain(
+// CHECK:<--- urKernelRetain(
 // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
 // CHECK-SAME: -> UR_RESULT_SUCCESS;
 //
-// CHECK:---> urCommandBufferAppendKernelLaunchExp(
+// CHECK:<--- urCommandBufferAppendKernelLaunchExp(
 // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
 //
-// CHECK:---> urKernelRelease(
+// CHECK:<--- urKernelRelease(
 // CHECK-SAME: .hKernel = [[KERNEL_HANDLE]]
 // CHECK-SAME: -> UR_RESULT_SUCCESS;
 

sycl/test-e2e/Graph/RecordReplay/memadvise.cpp

@@ -10,13 +10,13 @@
 // impact results but only performances, we verify
 // that a node is correctly added by checking UR function calls.
 
-// CHECK: urCommandBufferAppendUSMAdviseExp
+// CHECK: ---> urCommandBufferAppendUSMAdviseExp
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER:]]
 // CHECK-SAME: .pMemory = 0x[[#%x,PTR:]]
 // CHECK-SAME: .size = 400
 // CHECK-SAME: .pSyncPoint = {{.*}} (0x[[#%x,ADVISE_SYNC_POINT:]])
 
-// CHECK: urCommandBufferAppendKernelLaunchExp(
+// CHECK: ---> urCommandBufferAppendKernelLaunchExp(
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER]]
 // CHECK-SAME: .hKernel = 0x[[#%x,KERNEL:]]
 // CHECK-SAME: .workDim = 1

sycl/test-e2e/Graph/RecordReplay/prefetch.cpp

@@ -10,13 +10,13 @@
 // impact results but only performances, we verify
 // that a node is correctly added by checking UR function calls
 
-// CHECK: urCommandBufferAppendUSMPrefetchExp(
+// CHECK: ---> urCommandBufferAppendUSMPrefetchExp(
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER:]]
 // CHECK-SAME: .pMemory = 0x[[#%x,PTR:]]
 // CHECK-SAME: .size = 400
 // CHECK-SAME: .pSyncPoint = {{.*}} (0x[[#%x,PREFETCH_SYNC_POINT:]])
 
-// CHECK: urCommandBufferAppendKernelLaunchExp(
+// CHECK: ---> urCommandBufferAppendKernelLaunchExp(
 // CHECK-SAME: .hCommandBuffer = 0x[[#%x,COMMAND_BUFFER]]
 // CHECK-SAME: .hKernel = 0x[[#%x,KERNEL:]]
 // CHECK-SAME: .workDim = 1

sycl/test-e2e/KernelAndProgram/cache_env_vars.cpp

@@ -25,12 +25,12 @@
 // Some backends will call urProgramBuild and some will call
 // urProgramBuildExp depending on urProgramBuildExp support.
 
-// CHECK-BUILD-NOT: urProgramCreateWithBinary(
-// CHECK-BUILD: urProgramCreateWithIL(
-// CHECK-BUILD: urProgramBuild{{(Exp)?}}(
+// CHECK-BUILD-NOT: ---> urProgramCreateWithBinary(
+// CHECK-BUILD: ---> urProgramCreateWithIL(
+// CHECK-BUILD: ---> urProgramBuild{{(Exp)?}}(
 
-// CHECK-CACHE-NOT: urProgramCreateWithIL(
-// CHECK-CACHE: urProgramCreateWithBinary(
-// CHECK-CACHE: urProgramBuild{{(Exp)?}}(
+// CHECK-CACHE-NOT: ---> urProgramCreateWithIL(
+// CHECK-CACHE: ---> urProgramCreateWithBinary(
+// CHECK-CACHE: ---> urProgramBuild{{(Exp)?}}(
 
 #include "cache_env_vars.hpp"

sycl/test-e2e/KernelAndProgram/cache_env_vars_lin.cpp

@@ -24,12 +24,12 @@
 
 // Some backends will call urProgramBuild and some will call urProgramBuildExp depending on urProgramBuildExp support.
 
-// CHECK-BUILD-NOT: urProgramCreateWithBinary(
-// CHECK-BUILD: urProgramCreateWithIL(
-// CHECK-BUILD: urProgramBuild{{(Exp)?}}(
+// CHECK-BUILD-NOT: ---> urProgramCreateWithBinary(
+// CHECK-BUILD: ---> urProgramCreateWithIL(
+// CHECK-BUILD: ---> urProgramBuild{{(Exp)?}}(
 
-// CHECK-CACHE-NOT: urProgramCreateWithIL(
-// CHECK-CACHE: urProgramCreateWithBinary(
-// CHECK-CACHE: urProgramBuild{{(Exp)?}}(
+// CHECK-CACHE-NOT: ---> urProgramCreateWithIL(
+// CHECK-CACHE: ---> urProgramCreateWithBinary(
+// CHECK-CACHE: ---> urProgramBuild{{(Exp)?}}(
 
 #include "cache_env_vars.hpp"

sycl/test-e2e/KernelAndProgram/cache_env_vars_win.cpp

@@ -19,12 +19,12 @@
 // RUN: env SYCL_CACHE_PERSISTENT=1 AppData=%t/cache_dir SYCL_UR_TRACE=2 env -u SYCL_CACHE_DIR %t.out | FileCheck %s --check-prefixes=CHECK-BUILD
 // RUN: env SYCL_CACHE_PERSISTENT=1 AppData=%t/cache_dir SYCL_UR_TRACE=2 env -u SYCL_CACHE_DIR %t.out | FileCheck %s --check-prefixes=CHECK-CACHE
 
-// CHECK-BUILD-NOT: urProgramCreateWithBinary(
-// CHECK-BUILD: urProgramCreateWithIL(
-// CHECK-BUILD: urProgramBuild{{(Exp)?}}(
+// CHECK-BUILD-NOT: ---> urProgramCreateWithBinary(
+// CHECK-BUILD: ---> urProgramCreateWithIL(
+// CHECK-BUILD: ---> urProgramBuild{{(Exp)?}}(
 
-// CHECK-CACHE-NOT: urProgramCreateWithIL(
-// CHECK-CACHE: urProgramCreateWithBinary(
-// CHECK-CACHE: urProgramBuild{{(Exp)?}}(
+// CHECK-CACHE-NOT: ---> urProgramCreateWithIL(
+// CHECK-CACHE: ---> urProgramCreateWithBinary(
+// CHECK-CACHE: ---> urProgramBuild{{(Exp)?}}(
 
 #include "cache_env_vars.hpp"