[SYCL][sycl-post-link] Emit helpful warning message for undefined user functions in a fully linked device image (#17346)

In the following simple example (simplicity for sake of discussion),

```c++
#include <sycl/sycl.hpp>
SYCL_EXTERNAL int external_f1(int A, int B);
void hostf(...) {
    ...
    cgh.parallel_for<class Test>(range, [=](sycl::id<1> ID) {
      accC[ID] = external_f1(accA[ID], accB[ID]);
    });
  });
}
```

"external_f1" is a user function which is not defined in this module.
Linking this module stand-alone will successfully create an executable.
During run-time. there will be a run-time error saying that there is an
unresolved symbol. This PR addresses a user request to alert the user of
this behavior during compile-time by emitting a warning message.

Thanks

---------

Signed-off-by: Arvind Sudarsanam <[email protected]>

Author: asudarsa

llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

@@ -32,6 +32,7 @@
 #include "llvm/Support/Error.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/LineIterator.h"
+#include "llvm/Support/WithColor.h"
 #include "llvm/Transforms/IPO.h"
 #include "llvm/Transforms/IPO/GlobalDCE.h"
 #include "llvm/Transforms/IPO/Internalize.h"
@@ -298,6 +299,23 @@ void collectFunctionsAndGlobalVariablesToExtract(
   }
 }
 
+static bool isIntrinsicOrBuiltin(const Function &F) {
+  return F.isIntrinsic() || F.getName().starts_with("__") ||
+         isSpirvSyclBuiltin(F.getName()) || isESIMDBuiltin(F.getName());
+}
+
+// Checks for use of undefined user functions and emits a warning message.
+static void checkForCallsToUndefinedFunctions(const Module &M) {
+  if (AllowDeviceImageDependencies)
+    return;
+  for (const Function &F : M) {
+    if (!isIntrinsicOrBuiltin(F) && F.isDeclaration() && !F.use_empty())
+      WithColor::warning() << "Undefined function " << F.getName()
+                           << " found in " << M.getName()
+                           << ". This may result in runtime errors.\n";
+  }
+}
+
 // Check "spirv.ExecutionMode" named metadata in the module and remove nodes
 // that reference kernels that have dead prototypes or don't reference any
 // kernel at all (nullptr). Dead prototypes are removed as well.
@@ -382,6 +400,7 @@ ModuleDesc extractCallGraph(const ModuleDesc &MD,
   // GenXSPIRVWriterAdaptor pass that relies on this cleanup. This cleanup call
   // can be removed once that pass no longer depends on this cleanup.
   SplitM.cleanup();
+  checkForCallsToUndefinedFunctions(SplitM.getModule());
 
   return SplitM;
 }
@@ -1452,8 +1471,7 @@ bool canBeImportedFunction(const Function &F) {
   // in a header file.  Thus SYCL IR that is a declaration
   // will be considered as SYCL_EXTERNAL for the purposes of
   // this function.
-  if (F.isIntrinsic() || F.getName().starts_with("__") ||
-      isSpirvSyclBuiltin(F.getName()) || isESIMDBuiltin(F.getName()) ||
+  if (isIntrinsicOrBuiltin(F) ||
       (!F.isDeclaration() && !llvm::sycl::utils::isSYCLExternalFunction(&F)))
     return false;
 

sycl/test/warnings/undefined_functions.cpp

@@ -0,0 +1,25 @@
+// RUN: %clangxx %s -fsycl -fsycl-link 2>&1 | FileCheck %s --check-prefix=CHECK-WARNING
+// RUN: %clangxx %s -fsycl -fsycl-link -fsycl-allow-device-image-dependencies 2>&1 | FileCheck --allow-empty %s --check-prefix=CHECK-WARNING-DYNAMIC
+// This test is intended to check that we emit a helpful warning message for
+// undefined user functions in a fully linked device image after the
+// sycl-post-link stage of compilation.
+
+// CHECK-WARNING: warning: Undefined function _Z11external_f1ii found in
+// CHECK-WARNING-DYNAMIC-NOT: warning: Undefined function _Z11external_f1ii found in
+
+#include <sycl/sycl.hpp>
+
+SYCL_EXTERNAL int external_f1(int A, int B);
+
+void hostf(unsigned Size, sycl::buffer<int, 1> &bufA,
+           sycl::buffer<int, 1> &bufB, sycl::buffer<int, 1> &bufC) {
+  sycl::range<1> range{Size};
+  sycl::queue().submit([&](sycl::handler &cgh) {
+    auto accA = bufA.get_access<sycl::access::mode::read>(cgh);
+    auto accB = bufB.get_access<sycl::access::mode::read>(cgh);
+    auto accC = bufC.get_access<sycl::access::mode::write>(cgh);
+    cgh.parallel_for<class Test>(range, [=](sycl::id<1> ID) {
+      accC[ID] = external_f1(accA[ID], accB[ID]);
+    });
+  });
+}