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

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
hliao added reviewers: tra, rjmccall, yaxunl.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.

- Within `decltype`, expressions are only type-inspected. The restriction on 
CUDA calls should be relaxed.


Repository:
  rG LLVM Github Monorepo

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,17 @@
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
   bool IsAllowedCUDACall(const FunctionDecl *Caller,
                          const FunctionDecl *Callee) {
+    auto InDecltypeExpr = [this]() {
+      auto I =
+          std::find_if(ExprEvalContexts.rbegin(), ExprEvalContexts.rend(),
+                       [](const ExpressionEvaluationContextRecord &C) {
+                         return C.ExprContext ==
+                                ExpressionEvaluationContextRecord::EK_Decltype;
+                       });
+      return I != ExprEvalContexts.rend();
+    };
+    if (InDecltypeExpr())
+      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,17 @@
   /// semantically correct CUDA programs, but only if they're never codegen'ed.
   bool IsAllowedCUDACall(const FunctionDecl *Caller,
                          const FunctionDecl *Callee) {
+    auto InDecltypeExpr = [this]() {
+      auto I =
+          std::find_if(ExprEvalContexts.rbegin(), ExprEvalContexts.rend(),
+                       [](const ExpressionEvaluationContextRecord &C) {
+                         return C.ExprContext ==
+                                ExpressionEvaluationContextRecord::EK_Decltype;
+                       });
+      return I != ExprEvalContexts.rend();
+    };
+    if (InDecltypeExpr())
+      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