[SYCL] Test compiler diagnostics for unsupported printf usages (#5306)

Additionally, fix the diagnostic output for cases with an incorrect wrapper name and unify the punctuation style of the existing diagnostics. Signed-off-by: Artem Gindinson <[email protected]>
intel · Jan 15, 2022 · bf5d9d5 · bf5d9d5
1 parent 97e6d2d
commit bf5d9d5
Show file tree

Hide file tree

Showing 6 changed files with 213 additions and 2 deletions.
diff --git a/llvm/lib/SYCLLowerIR/MutatePrintfAddrspace.cpp b/llvm/lib/SYCLLowerIR/MutatePrintfAddrspace.cpp
@@ -210,7 +210,7 @@ size_t setFuncCallsOntoCASPrintf(Function *F, Function *CASPrintfFunc,
           "passing format strings directly into experimental::printf calls, "
           "avoiding indirection via wrapper function arguments.";
       if (!WrapperFunc->getName().contains("6oneapi12experimental6printf")) {
-        emitError(WrapperFunc, CI, BadWrapperErrorMsg);
+        emitError(F, CI, BadWrapperErrorMsg);
         return 0;
       }
       for (User *WrapperU : WrapperFunc->users()) {
@@ -239,7 +239,7 @@ size_t setFuncCallsOntoCASPrintf(Function *F, Function *CASPrintfFunc,
           F, CI,
           "Make sure each format string literal is "
           "known at compile time or use OpenCL constant address space literals "
-          "for device-side printf calls");
+          "for device-side printf calls.");
       return 0;
     }
   }

diff --git a/llvm/test/SYCLLowerIR/printf_addrspace/Inputs/experimental-printf-bad-inline-test.cpp b/llvm/test/SYCLLowerIR/printf_addrspace/Inputs/experimental-printf-bad-inline-test.cpp
@@ -0,0 +1,5 @@
+#include <CL/sycl.hpp>
+
+SYCL_EXTERNAL void custom_wrapper(const char *S) {
+  sycl::ext::oneapi::experimental::printf(S);
+}
diff --git a/llvm/test/SYCLLowerIR/printf_addrspace/Inputs/experimental-printf-compile-time-unknown.cpp b/llvm/test/SYCLLowerIR/printf_addrspace/Inputs/experimental-printf-compile-time-unknown.cpp
@@ -0,0 +1,18 @@
+#include <CL/sycl.hpp>
+
+using namespace sycl;
+
+int foo(int k) {
+  queue q;
+  buffer<int, 1> buf(k);
+  q.submit([&](handler &cgh) {
+    auto acc = buf.get_access(cgh);
+    cgh.single_task([=]() {
+      if (acc[0] == 0)
+        ext::oneapi::experimental::printf("String 0\n");
+      else
+        ext::oneapi::experimental::printf("String 1\n");
+    });
+  });
+  return 0;
+}
diff --git a/llvm/test/SYCLLowerIR/printf_addrspace/diagnose_bad_wrapper_inlining.ll b/llvm/test/SYCLLowerIR/printf_addrspace/diagnose_bad_wrapper_inlining.ll
@@ -0,0 +1,47 @@
+;; This tests diagnostic emission whenever printf's format string argument
+;; is taken from the argument list of a user wrapper function, but then the
+;; user-provided wrapper does not get inlined (while the library-provided
+;; experimental::printf wrapper does).
+
+;; The IR for test purposes is based on the following source/compilation (custom
+;; build of SYCL Clang with SYCLMutatePrintfAddrspacePass turned off):
+;; clang++ -fsycl -fsycl-device-only Inputs/experimental-printf-bad-inline-test.cpp -S -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ -O2
+
+; RUN: not opt < %s --SYCLMutatePrintfAddrspace -S --enable-new-pm=0 2>&1 | FileCheck %s
+; RUN: not opt < %s --passes=SYCLMutatePrintfAddrspace -S 2>&1 | FileCheck %s
+; CHECK: error: experimental::printf requires format string to reside in constant address space. The compiler wasn't able to automatically convert your format string into constant address space when processing builtin _Z18__spirv_ocl_printf{{.*}} called in function {{.*}}custom_wrapper{{.*}}.
+; CHECK-NEXT: Consider simplifying the code by passing format strings directly into experimental::printf calls, avoiding indirection via wrapper function arguments.
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"class.cl::sycl::id" = type { %"class.cl::sycl::detail::array" }
+%"class.cl::sycl::detail::array" = type { [1 x i64] }
+
+; Function Attrs: convergent mustprogress norecurse
+define dso_local spir_func void @_Z14custom_wrapperPKc(i8 addrspace(4)* %S) local_unnamed_addr #0 {
+entry:
+  %call.i = tail call spir_func i32 @_Z18__spirv_ocl_printfIJEEiPKcDpT_(i8 addrspace(4)* %S) #3
+  ret void
+}
+
+; Function Attrs: convergent
+declare dso_local spir_func i32 @_Z18__spirv_ocl_printfIJEEiPKcDpT_(i8 addrspace(4)*) local_unnamed_addr #1
+
+attributes #0 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf.cpp" }
+attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf-bad-inline-test.cpp" "uniform-work-group-size"="true" }
+attributes #3 = { convergent }
+
+!llvm.module.flags = !{!0, !1}
+!opencl.spir.version = !{!2}
+!spirv.Source = !{!3}
+!llvm.ident = !{!4}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"frame-pointer", i32 2}
+!2 = !{i32 1, i32 2}
+!3 = !{i32 4, i32 100000}
+!4 = !{!"clang version 14.0.0"}
+!5 = !{i32 -1, i32 -1, i32 -1, i32 -1}
+!6 = !{i1 false, i1 true, i1 true, i1 false}
diff --git a/llvm/test/SYCLLowerIR/printf_addrspace/diagnose_compile_time_unknown.ll b/llvm/test/SYCLLowerIR/printf_addrspace/diagnose_compile_time_unknown.ll
@@ -0,0 +1,63 @@
+;; This tests diagnostic emission upon compile-time-unknown format string.
+;; In this instance, a select instruction is given as the first argument
+;; to __spirv_ocl_printf - as a result, moving the function into the constant AS
+;; becomes impossible.
+
+;; The IR is based on the following source/compilation (custom
+;; build of SYCL Clang with SYCLMutatePrintfAddrspacePass turned off):
+;; clang++ -fsycl -fsycl-device-only Inputs/experimental-printf-compile-time-unknown.cpp -S -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
+
+; RUN: not opt < %s --SYCLMutatePrintfAddrspace -S --enable-new-pm=0 2>&1 | FileCheck %s
+; RUN: not opt < %s --passes=SYCLMutatePrintfAddrspace -S 2>&1 | FileCheck %s
+; CHECK: error: experimental::printf requires format string to reside in constant address space. The compiler wasn't able to automatically convert your format string into constant address space when processing builtin _Z18__spirv_ocl_printf{{.*}} called in function {{.*}}foo{{.*}}.
+; CHECK-NEXT: Make sure each format string literal is known at compile time or use OpenCL constant address space literals for device-side printf calls.
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"class.cl::sycl::id" = type { %"class.cl::sycl::detail::array" }
+%"class.cl::sycl::detail::array" = type { [1 x i64] }
+
+$_ZTSZZ3fooiENKUlRN2cl4sycl7handlerEE_clES2_EUlvE_ = comdat any
+
+@.str = private unnamed_addr addrspace(1) constant [10 x i8] c"String 0\0A\00", align 1
+@.str.1 = private unnamed_addr addrspace(1) constant [10 x i8] c"String 1\0A\00", align 1
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @_ZTSZZ3fooiENKUlRN2cl4sycl7handlerEE_clES2_EUlvE_(i32 addrspace(1)* %_arg_, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !5 !sycl_kernel_omit_args !6 {
+entry:
+  %0 = getelementptr inbounds %"class.cl::sycl::id", %"class.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+  %1 = addrspacecast i64* %0 to i64 addrspace(4)*
+  %2 = load i64, i64 addrspace(4)* %1, align 8
+  %add.ptr.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_, i64 %2
+  %arrayidx.ascast.i.i = addrspacecast i32 addrspace(1)* %add.ptr.i to i32 addrspace(4)*
+  %3 = load i32, i32 addrspace(4)* %arrayidx.ascast.i.i, align 4, !tbaa !7
+  %cmp.i = icmp eq i32 %3, 0
+  %..i = select i1 %cmp.i, i8 addrspace(4)* getelementptr inbounds ([10 x i8], [10 x i8] addrspace(4)* addrspacecast ([10 x i8] addrspace(1)* @.str to [10 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* getelementptr inbounds ([10 x i8], [10 x i8] addrspace(4)* addrspacecast ([10 x i8] addrspace(1)* @.str.1 to [10 x i8] addrspace(4)*), i64 0, i64 0)
+  %call.i.i = tail call spir_func i32 @_Z18__spirv_ocl_printfIJEEiPKcDpT_(i8 addrspace(4)* %..i) #2
+  ret void
+}
+
+; Function Attrs: convergent
+declare dso_local spir_func i32 @_Z18__spirv_ocl_printfIJEEiPKcDpT_(i8 addrspace(4)*) local_unnamed_addr #1
+
+attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf.cpp" "uniform-work-group-size"="true" }
+attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { convergent }
+
+!llvm.module.flags = !{!0, !1}
+!opencl.spir.version = !{!2}
+!spirv.Source = !{!3}
+!llvm.ident = !{!4}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"frame-pointer", i32 2}
+!2 = !{i32 1, i32 2}
+!3 = !{i32 4, i32 100000}
+!4 = !{!"clang version 14.0.0"}
+!5 = !{i32 -1, i32 -1, i32 -1, i32 -1}
+!6 = !{i1 false, i1 true, i1 true, i1 false}
+!7 = !{!8, !8, i64 0}
+!8 = !{!"int", !9, i64 0}
+!9 = !{!"omnipotent char", !10, i64 0}
+!10 = !{!"Simple C++ TBAA"}
diff --git a/llvm/test/SYCLLowerIR/printf_addrspace/diagnose_extra_user_wrapper_no_opt.ll b/llvm/test/SYCLLowerIR/printf_addrspace/diagnose_extra_user_wrapper_no_opt.ll
@@ -0,0 +1,78 @@
+;; This tests diagnostic emission whenever printf's format string argument
+;; is taken from the argument list of a user wrapper function and no
+;; inlining is performed prior to the address space mutation pass.
+
+;; The IR for test purposes is based on the following source/compilation (custom
+;; build of SYCL Clang with SYCLMutatePrintfAddrspacePass turned off):
+;; clang++ -fsycl -fsycl-device-only Inputs/experimental-printf-bad-inline-test.cpp -S -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__ -O0
+
+; RUN: not opt < %s --SYCLMutatePrintfAddrspace -S --enable-new-pm=0 2>&1 | FileCheck %s
+; RUN: not opt < %s --passes=SYCLMutatePrintfAddrspace -S 2>&1 | FileCheck %s
+; CHECK: error: experimental::printf requires format string to reside in constant address space. The compiler wasn't able to automatically convert your format string into constant address space when processing builtin _ZN2cl4sycl3ext6oneapi12experimental6printf{{.*}} called in function {{.*}}custom_wrapper{{.*}}.
+; CHECK-NEXT: Consider simplifying the code by passing format strings directly into experimental::printf calls, avoiding indirection via wrapper function arguments.
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" }
+%"class.cl::sycl::detail::array" = type { [1 x i64] }
+%"class.cl::sycl::id" = type { %"class.cl::sycl::detail::array" }
+%class.anon = type { %"class.cl::sycl::accessor" }
+%"class.cl::sycl::accessor" = type { %"class.cl::sycl::detail::AccessorImplDevice" }
+%"class.cl::sycl::detail::AccessorImplDevice" = type { %"class.cl::sycl::id", %"class.cl::sycl::range", %"class.cl::sycl::range" }
+%"class.cl::sycl::detail::accessor_common" = type { i8 }
+
+$_ZN2cl4sycl3ext6oneapi12experimental6printfIcJEEEiPKT_DpT0_ = comdat any
+
+; Function Attrs: convergent mustprogress noinline norecurse optnone
+define dso_local spir_func void @_Z14custom_wrapperPKc(i8 addrspace(4)* %S) #0 {
+entry:
+  %S.addr = alloca i8 addrspace(4)*, align 8
+  %S.addr.ascast = addrspacecast i8 addrspace(4)** %S.addr to i8 addrspace(4)* addrspace(4)*
+  store i8 addrspace(4)* %S, i8 addrspace(4)* addrspace(4)* %S.addr.ascast, align 8
+  %0 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %S.addr.ascast, align 8
+  %call = call spir_func i32 @_ZN2cl4sycl3ext6oneapi12experimental6printfIcJEEEiPKT_DpT0_(i8 addrspace(4)* %0) #9
+  ret void
+}
+
+; Function Attrs: convergent mustprogress noinline norecurse optnone
+define linkonce_odr dso_local spir_func i32 @_ZN2cl4sycl3ext6oneapi12experimental6printfIcJEEEiPKT_DpT0_(i8 addrspace(4)* %__format) #1 comdat {
+entry:
+  %retval = alloca i32, align 4
+  %__format.addr = alloca i8 addrspace(4)*, align 8
+  %retval.ascast = addrspacecast i32* %retval to i32 addrspace(4)*
+  %__format.addr.ascast = addrspacecast i8 addrspace(4)** %__format.addr to i8 addrspace(4)* addrspace(4)*
+  store i8 addrspace(4)* %__format, i8 addrspace(4)* addrspace(4)* %__format.addr.ascast, align 8
+  %0 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %__format.addr.ascast, align 8
+  %call = call spir_func i32 @_Z18__spirv_ocl_printfIJEEiPKcDpT_(i8 addrspace(4)* %0) #9
+  ret i32 %call
+}
+
+; Function Attrs: convergent
+declare dso_local spir_func i32 @_Z18__spirv_ocl_printfIJEEiPKcDpT_(i8 addrspace(4)*) #2
+
+attributes #0 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf-bad-inline-test.cpp" }
+attributes #1 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf-bad-inline-test.cpp" "uniform-work-group-size"="true" }
+attributes #4 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #5 = { argmemonly nofree nounwind willreturn }
+attributes #6 = { argmemonly nofree nounwind willreturn writeonly }
+attributes #7 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #8 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #9 = { convergent }
+
+!llvm.module.flags = !{!0, !1}
+!opencl.spir.version = !{!2}
+!spirv.Source = !{!3}
+!llvm.ident = !{!4}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"frame-pointer", i32 2}
+!2 = !{i32 1, i32 2}
+!3 = !{i32 4, i32 100000}
+!4 = !{!"clang version 14.0.0"}
+!5 = !{i32 -1, i32 -1, i32 -1, i32 -1}
+!6 = distinct !{!6, !7, !8}
+!7 = !{!"llvm.loop.mustprogress"}
+!8 = !{!"llvm.loop.unroll.enable"}