r352488 - [CUDA][HIP] Do not diagnose use of _Float16

[Yaxun Liu via cfe-commits]

Author: yaxunl
Date: Tue Jan 29 05:20:23 2019
New Revision: 352488

URL: http://llvm.org/viewvc/llvm-project?rev=352488&view=rev
Log:
[CUDA][HIP] Do not diagnose use of _Float16

r352221 caused regressions in CUDA/HIP since device function may use _Float16 whereas host does not support it.
In this case host compilation should not diagnose usage of _Float16 in device functions or variables.

For now just do not diagnose _Float16 for CUDA/HIP. In the future we should have more precise check.

Differential Revision: https://reviews.llvm.org/D57369

Added:
    cfe/trunk/test/SemaCUDA/float16.cu
Modified:
    cfe/trunk/lib/Lex/LiteralSupport.cpp
    cfe/trunk/lib/Sema/SemaType.cpp

Modified: cfe/trunk/lib/Lex/LiteralSupport.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Lex/LiteralSupport.cpp?rev=352488&r1=352487&r2=352488&view=diff
==============================================================================
--- cfe/trunk/lib/Lex/LiteralSupport.cpp (original)
+++ cfe/trunk/lib/Lex/LiteralSupport.cpp Tue Jan 29 05:20:23 2019
@@ -616,8 +616,11 @@ NumericLiteralParser::NumericLiteralPars
       if (isHalf || isFloat || isLong || isFloat128)
         break; // HF, FF, LF, QF invalid.
 
-      if (PP.getTargetInfo().hasFloat16Type() && s + 2 < ThisTokEnd &&
-          s[1] == '1' && s[2] == '6') {
+      // CUDA host and device may have different _Float16 support, therefore
+      // allows f16 literals to avoid false alarm.
+      // ToDo: more precise check for CUDA.
+      if ((PP.getTargetInfo().hasFloat16Type() || PP.getLangOpts().CUDA) &&
+          s + 2 < ThisTokEnd && s[1] == '1' && s[2] == '6') {
         s += 2; // success, eat up 2 characters.
         isFloat16 = true;
         continue;

Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=352488&r1=352487&r2=352488&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Tue Jan 29 05:20:23 2019
@@ -1442,7 +1442,10 @@ static QualType ConvertDeclSpecToType(Ty
       Result = Context.Int128Ty;
     break;
   case DeclSpec::TST_float16:
-    if (!S.Context.getTargetInfo().hasFloat16Type())
+    // CUDA host and device may have different _Float16 support, therefore
+    // do not diagnose _Float16 usage to avoid false alarm.
+    // ToDo: more precise diagnostics for CUDA.
+    if (!S.Context.getTargetInfo().hasFloat16Type() && !S.getLangOpts().CUDA)
       S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
         << "_Float16";
     Result = Context.Float16Ty;

Added: cfe/trunk/test/SemaCUDA/float16.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/float16.cu?rev=352488&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/float16.cu (added)
+++ cfe/trunk/test/SemaCUDA/float16.cu Tue Jan 29 05:20:23 2019
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify %s
+// expected-no-diagnostics
+#include "Inputs/cuda.h"
+
+__device__ void f(_Float16 x);
+
+__device__ _Float16 x = 1.0f16;