[CUDA] CUDA has no device-side library builtins.

[tra]

Summary

We should (almost) never consider a device-side declaration to match a
builtin. If we do, the un-inlined device-side functions provided by
CUDA headers that ship with clang may be ignored. We may end up emitting
as a call to a llvm intrinsic which would typically be lowered as
an external library call. This results in a back-end error because NVPTX
back-end does not support it.

[jlebar]

How does this affect e.g. calling memcpy()? There isn't a standard library implementation of this on nvptx, but we do want calls to memcpy() to be lowered to llvm.memcpy so that they can be optimized.

[tra]

In D42319#983377, @jlebar wrote:

How does this affect e.g. calling memcpy()? There isn't a standard library implementation of this on nvptx, but we do want calls to memcpy() to be lowered to llvm.memcpy so that they can be optimized.

We implement memcpy as a call to __builtin_memcpy() which gets code-gen-ed as usual. NVPTX also lowers all memcpy/memset/memmove as loads/stores, so those don't need external library. This behavior is not affected by this patch.

This patch's goal is to prevent clang codegen-ing its idea of the library builtin function while ignoring the implementation we've provided in the headers for device side.

Original issue I had was triggered by code roughly similar to this:

extern "C" __device__ int logf(float a) { return __nv_logf(a); }
__global__ void kernel() { logf(0.0f); }

In the AST, the kernel was calling the logf functions above However, when clang generated code, it considered that logf is a library builtin with known semantics and happily codegen'ed a call to @llvm.log.f32, which NVPTX back-end has no way to lower. The patch adds a safety net in clang so it does not generate code for builtins which we have disabled (or can't handle) in NVPTX.

[jlebar]

Got it, thanks for the explanation.

clang/lib/AST/Decl.cpp

if (getStorageClass() == SC_Static)
     return 0;
 
   // OpenCL v1.2 s6.9.f - The library functions defined in
   // the C99 standard headers are not available.
   if (Context.getLangOpts().OpenCL &&
       Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID))
     return 0;
 
+  // CUDA does not have device-side standard library. printf and malloc are the
+  // only special cases that are supported by device-side runtime.
+  if (Context.getLangOpts().CUDA && hasAttr<CUDADeviceAttr>() &&
+      !hasAttr<CUDAHostAttr>() &&
+      !(BuiltinID == Builtin::BIprintf || BuiltinID == Builtin::BImalloc))
+    return 0;
+
   return BuiltinID;
 }
 
 /// getNumParams - Return the number of parameters this function must have
 /// based on its FunctionType.  This is the length of the ParamInfo array
 /// after it has been created.
 unsigned FunctionDecl::getNumParams() const {
   const auto *FPT = getType()->getAs<FunctionProtoType>();

clang/test/CodeGenCUDA/library-builtin.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | \
+// RUN:  FileCheck --check-prefixes=HOST,BOTH %s
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \
+// RUN:   -emit-llvm -o - %s | FileCheck %s --check-prefixes=DEVICE,BOTH
+
+// BOTH-LABEL: define float @logf(float
+
+// logf() should be calling itself recursively as we don't have any standard
+// library on device side.
+// DEVICE: call float @logf(float
+extern "C" __attribute__((device)) float logf(float __x) { return logf(__x); }
+
+// NOTE: this case is to illustrate the expected differences in behavior between
+// the host and device. In general we do not mess with host-side standard
+// library.
+//
+// Host is assumed to have standerd library, so logf() calls LLVM intrinsic.
+// HOST: call float @llvm.log.f32(float
+extern "C" float logf(float __x) { return logf(__x); }