[PATCH] D61458: [hip] Relax CUDA call restriction within `decltype` context.

[Michael Liao via Phabricator via cfe-commits]

hliao updated this revision to Diff 197860.
hliao added a comment.

simplify the logic using `llvm::any_of`.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D61458/new/

https://reviews.llvm.org/D61458

Files:
  clang/include/clang/Sema/Sema.h
  clang/test/CodeGenCUDA/function-overload.cu


Index: clang/test/CodeGenCUDA/function-overload.cu
===================================================================
--- clang/test/CodeGenCUDA/function-overload.cu
+++ clang/test/CodeGenCUDA/function-overload.cu
@@ -8,6 +8,8 @@
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o 
- %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+// RUN: %clang_cc1 -std=c++11 -DCHECK_DECLTYPE -triple amdgcn -fcuda-is-device 
-emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-DECLTYPE %s
 
 #include "Inputs/cuda.h"
 
@@ -53,3 +55,14 @@
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
 // CHECK-BOTH: store i32 32,
 // CHECK-BOTH: ret void
+
+#if defined(CHECK_DECLTYPE)
+int foo(float);
+// CHECK-DECLTYPE-LABEL: @_Z3barf
+// CHECK-DECLTYPE: fptosi
+// CHECK-DECLTYPE: sitofp
+__device__ float bar(float x) {
+  decltype(foo(x)) y = x;
+  return y + 3.f;
+}
+#endif
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -10406,6 +10406,12 @@
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
   bool IsAllowedCUDACall(const FunctionDecl *Caller,
                          const FunctionDecl *Callee) {
+    if (llvm::any_of(ExprEvalContexts,
+                     [](const ExpressionEvaluationContextRecord &C) {
+                       return C.ExprContext ==
+                              ExpressionEvaluationContextRecord::EK_Decltype;
+                     }))
+      return true;
     return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
   }
 



Index: clang/test/CodeGenCUDA/function-overload.cu
===================================================================
--- clang/test/CodeGenCUDA/function-overload.cu
+++ clang/test/CodeGenCUDA/function-overload.cu
@@ -8,6 +8,8 @@
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+// RUN: %clang_cc1 -std=c++11 -DCHECK_DECLTYPE -triple amdgcn -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-DECLTYPE %s
 
 #include "Inputs/cuda.h"
 
@@ -53,3 +55,14 @@
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
 // CHECK-BOTH: store i32 32,
 // CHECK-BOTH: ret void
+
+#if defined(CHECK_DECLTYPE)
+int foo(float);
+// CHECK-DECLTYPE-LABEL: @_Z3barf
+// CHECK-DECLTYPE: fptosi
+// CHECK-DECLTYPE: sitofp
+__device__ float bar(float x) {
+  decltype(foo(x)) y = x;
+  return y + 3.f;
+}
+#endif
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -10406,6 +10406,12 @@
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
   bool IsAllowedCUDACall(const FunctionDecl *Caller,
                          const FunctionDecl *Callee) {
+    if (llvm::any_of(ExprEvalContexts,
+                     [](const ExpressionEvaluationContextRecord &C) {
+                       return C.ExprContext ==
+                              ExpressionEvaluationContextRecord::EK_Decltype;
+                     }))
+      return true;
     return IdentifyCUDAPreference(Caller, Callee) != CFP_Never;
   }
 

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits