From d4dffe409b9bf7bc7e3c54937548c76b7c0a9166 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Mon, 9 Aug 2021 10:03:24 +0100 Subject: [PATCH 1/3] [SYCL] Add support for NVPTX device printf Use `::printf` when not compiling for `__SPIR__`, this allows the use of `EmitNVPTXDevicePrintfCallExpr` which packs the var args and dispatches to CUDA's `vprintf`. Fixes #1154 --- clang/lib/CodeGen/CGGPUBuiltin.cpp | 15 +++++++++++++-- clang/lib/Sema/SemaSYCL.cpp | 2 ++ .../sycl/ext/oneapi/experimental/builtins.hpp | 4 ++-- 3 files changed, 17 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp b/clang/lib/CodeGen/CGGPUBuiltin.cpp index f860623e2bc37..205a9dd49a7af 100644 --- a/clang/lib/CodeGen/CGGPUBuiltin.cpp +++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp @@ -118,8 +118,19 @@ CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E, // Invoke vprintf and return. llvm::Function* VprintfFunc = GetVprintfDeclaration(CGM.getModule()); - return RValue::get(Builder.CreateCall( - VprintfFunc, {Args[0].getRValue(*this).getScalarVal(), BufferPtr})); + auto FormatSpecifier = Args[0].getRValue(*this).getScalarVal(); + // Check if the format specifier is in the constant address space, vprintf is + // oblivious to address spaces, so it would have to be casted away. + if (Args[0] + .getRValue(*this) + .getScalarVal() + ->getType() + ->getPointerAddressSpace() == 4) + FormatSpecifier = Builder.CreateAddrSpaceCast( + FormatSpecifier, llvm::Type::getInt8PtrTy(Ctx)); + + return RValue::get( + Builder.CreateCall(VprintfFunc, {FormatSpecifier, BufferPtr})); } RValue diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0e196da8325dd..41fd64edb9b69 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -420,6 +420,8 @@ static bool IsSyclMathFunc(unsigned BuiltinID) { bool Sema::isKnownGoodSYCLDecl(const Decl *D) { if (const FunctionDecl *FD = dyn_cast(D)) { const IdentifierInfo *II = FD->getIdentifier(); + if (FD->getBuiltinID() == Builtin::BIprintf) + return true; const DeclContext *DC = FD->getDeclContext(); if (II && II->isStr("__spirv_ocl_printf") && !FD->isDefined() && diff --git a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp index ee93897b1ef65..7108e189f51da 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp @@ -61,11 +61,11 @@ namespace experimental { // template int printf(const __SYCL_CONSTANT_AS char *__format, Args... args) { -#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) return __spirv_ocl_printf(__format, args...); #else return ::printf(__format, args...); -#endif +#endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) } } // namespace experimental From f7bbcc6d10832ec45fec34d9a8881a782ea264d1 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Fri, 13 Aug 2021 09:28:31 +0100 Subject: [PATCH 2/3] [SYCL] Add NVPTX printf test and printf support in mock SYCL header --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 17 +++++++++++++++++ clang/test/CodeGenSYCL/nvptx-printf.cpp | 18 ++++++++++++++++++ 2 files changed, 35 insertions(+) create mode 100644 clang/test/CodeGenSYCL/nvptx-printf.cpp diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 957b706926d1a..5fed15a1a6acc 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -2,6 +2,8 @@ #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) +extern "C" int printf(const char* fmt, ...); + // Dummy runtime classes to model SYCL API. inline namespace cl { namespace sycl { @@ -310,6 +312,21 @@ class spec_constant { return get(); } }; + +#ifdef __SYCL_DEVICE_ONLY__ +#define __SYCL_CONSTANT_AS __attribute__((opencl_constant)) +#else +#define __SYCL_CONSTANT_AS +#endif +template +int printf(const __SYCL_CONSTANT_AS char *__format, Args... args) { +#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) + return __spirv_ocl_printf(__format, args...); +#else + return ::printf(__format, args...); +#endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +} + } // namespace experimental } // namespace oneapi } // namespace ext diff --git a/clang/test/CodeGenSYCL/nvptx-printf.cpp b/clang/test/CodeGenSYCL/nvptx-printf.cpp new file mode 100644 index 0000000000000..9be018f90b870 --- /dev/null +++ b/clang/test/CodeGenSYCL/nvptx-printf.cpp @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-nvidia-cuda-sycldevice -std=c++11 -S -emit-llvm -x c++ %s -o - | FileCheck %s + +#include "Inputs/sycl.hpp" + +#ifdef __SYCL_DEVICE_ONLY__ +#define CONSTANT __attribute__((opencl_constant)) +#else +#define CONSTANT +#endif + +static const CONSTANT char format_2[] = "Hello! %d %f\n"; + +int main() { + // Make sure that device printf is dispatched to CUDA's vprintf syscall. + // CHECK: alloca %printf_args + // CHECK: call i32 @vprintf + cl::sycl::kernel_single_task([]() { cl::sycl::ext::oneapi::experimental::printf(format_2, 123, 1.23); }); +} From c8b0ea25ade93ee7f9e83159600bc206fea19938 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Tue, 17 Aug 2021 13:08:02 +0300 Subject: [PATCH 3/3] Update clang/test/CodeGenSYCL/nvptx-printf.cpp --- clang/test/CodeGenSYCL/nvptx-printf.cpp | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/clang/test/CodeGenSYCL/nvptx-printf.cpp b/clang/test/CodeGenSYCL/nvptx-printf.cpp index 9be018f90b870..8246c0fe1ab81 100644 --- a/clang/test/CodeGenSYCL/nvptx-printf.cpp +++ b/clang/test/CodeGenSYCL/nvptx-printf.cpp @@ -2,13 +2,7 @@ #include "Inputs/sycl.hpp" -#ifdef __SYCL_DEVICE_ONLY__ -#define CONSTANT __attribute__((opencl_constant)) -#else -#define CONSTANT -#endif - -static const CONSTANT char format_2[] = "Hello! %d %f\n"; +static const __SYCL_CONSTANT_AS char format_2[] = "Hello! %d %f\n"; int main() { // Make sure that device printf is dispatched to CUDA's vprintf syscall.