[PATCH] D67509: [CUDA][HIP] Fix hostness of defaulted constructor

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl updated this revision to Diff 220838.
yaxunl retitled this revision from "[CUDA][HIP] Diagnose defaulted constructor 
only if it is used" to "[CUDA][HIP] Fix hostness of defaulted constructor".
yaxunl edited the summary of this revision.
yaxunl added a comment.

Posts a new fix for this issue, where the defaulted constructor definition 
follows the hostness of the original declaration in the class. Also fix the 
issue when defaulted ctor has explicit host device attribs.


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

https://reviews.llvm.org/D67509

Files:
  lib/Sema/SemaCUDA.cpp
  test/SemaCUDA/default-ctor.cu


Index: test/SemaCUDA/default-ctor.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/default-ctor.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:            -fcuda-is-device -verify -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+struct In { In() = default; };
+struct InD { __device__ InD() = default; };
+struct InH { __host__ InH() = default; };
+struct InHD { __host__ __device__ InHD() = default; };
+
+struct Out { Out(); };
+struct OutD { __device__ OutD(); };
+struct OutH { __host__ OutH(); };
+struct OutHD { __host__ __device__ OutHD(); };
+
+Out::Out() = default;
+__device__ OutD::OutD() = default;
+__host__ OutH::OutH() = default;
+__host__ __device__ OutHD::OutHD() = default;
+
+__device__ void fd() {
+  In in;
+  InD ind;
+  InH inh; // expected-error{{no matching constructor for initialization of 
'InH'}}
+  InHD inhd;
+  Out out; // expected-error{{no matching constructor for initialization of 
'Out'}}
+  OutD outd;
+  OutH outh; // expected-error{{no matching constructor for initialization of 
'OutH'}}
+  OutHD outhd;
+}
+
+__host__ void fh() {
+  In in;
+  InD ind; // expected-error{{no matching constructor for initialization of 
'InD'}}
+  InH inh;
+  InHD inhd;
+  Out out;
+  OutD outd; // expected-error{{no matching constructor for initialization of 
'OutD'}}
+  OutH outh;
+  OutHD outhd;
+}
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -267,6 +267,12 @@
                                                    CXXMethodDecl *MemberDecl,
                                                    bool ConstRHS,
                                                    bool Diagnose) {
+  bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
+  bool hasAttr = MemberDecl->hasAttr<CUDADeviceAttr>() ||
+                 MemberDecl->hasAttr<CUDAHostAttr>();
+  if (!InClass || hasAttr)
+    return false;
+
   llvm::Optional<CUDAFunctionTarget> InferredTarget;
 
   // We're going to invoke special member lookup; mark that these special



Index: test/SemaCUDA/default-ctor.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/default-ctor.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:            -fcuda-is-device -verify -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+struct In { In() = default; };
+struct InD { __device__ InD() = default; };
+struct InH { __host__ InH() = default; };
+struct InHD { __host__ __device__ InHD() = default; };
+
+struct Out { Out(); };
+struct OutD { __device__ OutD(); };
+struct OutH { __host__ OutH(); };
+struct OutHD { __host__ __device__ OutHD(); };
+
+Out::Out() = default;
+__device__ OutD::OutD() = default;
+__host__ OutH::OutH() = default;
+__host__ __device__ OutHD::OutHD() = default;
+
+__device__ void fd() {
+  In in;
+  InD ind;
+  InH inh; // expected-error{{no matching constructor for initialization of 'InH'}}
+  InHD inhd;
+  Out out; // expected-error{{no matching constructor for initialization of 'Out'}}
+  OutD outd;
+  OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}}
+  OutHD outhd;
+}
+
+__host__ void fh() {
+  In in;
+  InD ind; // expected-error{{no matching constructor for initialization of 'InD'}}
+  InH inh;
+  InHD inhd;
+  Out out;
+  OutD outd; // expected-error{{no matching constructor for initialization of 'OutD'}}
+  OutH outh;
+  OutHD outhd;
+}
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -267,6 +267,12 @@
                                                    CXXMethodDecl *MemberDecl,
                                                    bool ConstRHS,
                                                    bool Diagnose) {
+  bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
+  bool hasAttr = MemberDecl->hasAttr<CUDADeviceAttr>() ||
+                 MemberDecl->hasAttr<CUDAHostAttr>();
+  if (!InClass || hasAttr)
+    return false;
+
   llvm::Optional<CUDAFunctionTarget> InferredTarget;
 
   // We're going to invoke special member lookup; mark that these special

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