[clang] 9c2cfaa - [Clang][OpenMP] Allow `f16` literal suffix when compiling OpenMP target offloading for NVPTX

[Shilei Tian via cfe-commits]

Author: Shilei Tian
Date: 2023-01-19T22:24:38-05:00
New Revision: 9c2cfaaada66f323d10e25ac961ce1524d7b5f1f

URL: https://github.com/llvm/llvm-project/commit/9c2cfaaada66f323d10e25ac961ce1524d7b5f1f
DIFF: https://github.com/llvm/llvm-project/commit/9c2cfaaada66f323d10e25ac961ce1524d7b5f1f.diff

LOG: [Clang][OpenMP] Allow `f16` literal suffix when compiling OpenMP target offloading for NVPTX

Fix #58087.

Reviewed By: jhuber6

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

Added: 
    clang/test/OpenMP/float16_sema.cpp

Modified: 
    clang/lib/Lex/LiteralSupport.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Lex/LiteralSupport.cpp b/clang/lib/Lex/LiteralSupport.cpp
index fb2b14601a6ae..421a853360430 100644
--- a/clang/lib/Lex/LiteralSupport.cpp
+++ b/clang/lib/Lex/LiteralSupport.cpp
@@ -943,9 +943,13 @@ NumericLiteralParser::NumericLiteralParser(StringRef TokSpelling,
 
       // CUDA host and device may have 
diff erent _Float16 support, therefore
       // allows f16 literals to avoid false alarm.
+      // When we compile for OpenMP target offloading on NVPTX, f16 suffix
+      // should also be supported.
       // ToDo: more precise check for CUDA.
-      if ((Target.hasFloat16Type() || LangOpts.CUDA) && s + 2 < ThisTokEnd &&
-          s[1] == '1' && s[2] == '6') {
+      // TODO: AMDGPU might also support it in the future.
+      if ((Target.hasFloat16Type() || LangOpts.CUDA ||
+           (LangOpts.OpenMPIsDevice && Target.getTriple().isNVPTX())) &&
+          s + 2 < ThisTokEnd && s[1] == '1' && s[2] == '6') {
         s += 2; // success, eat up 2 characters.
         isFloat16 = true;
         continue;

diff  --git a/clang/test/OpenMP/float16_sema.cpp b/clang/test/OpenMP/float16_sema.cpp
new file mode 100644
index 0000000000000..4b9f3399eb30d
--- /dev/null
+++ b/clang/test/OpenMP/float16_sema.cpp
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -fsyntax-only -x c++ -triple x86_64-unknown-linux-gnu -fopenmp -fopenmp-targets=nvptx64 -verify %s
+// expected-no-diagnostics
+
+int foo() {
+#pragma omp target
+  {
+    __fp16 a = -1.0f16;
+  }
+  return 0;
+}