Author: tra Date: Tue Sep 13 17:16:30 2016 New Revision: 281406 URL: http://llvm.org/viewvc/llvm-project?rev=281406&view=rev Log: [CUDA] Do not merge CUDA target attributes.
CUDA target attributes are used for function overloading and must not be merged. This fixes a bug where attributes were inherited during function template specialization in CUDA and made it impossible for specialized function to provide its own target attributes. Differential Revision: https://reviews.llvm.org/D24522 Added: cfe/trunk/test/SemaCUDA/target_attr_inheritance.cu Modified: cfe/trunk/lib/Sema/SemaDecl.cpp cfe/trunk/test/SemaCUDA/function-overload.cu Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=281406&r1=281405&r2=281406&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Sep 13 17:16:30 2016 @@ -2290,7 +2290,13 @@ static bool mergeDeclAttribute(Sema &S, NewAttr = S.mergeAlwaysInlineAttr(D, AA->getRange(), &S.Context.Idents.get(AA->getSpelling()), AttrSpellingListIndex); - else if (const auto *MA = dyn_cast<MinSizeAttr>(Attr)) + else if (S.getLangOpts().CUDA && isa<FunctionDecl>(D) && + (isa<CUDAHostAttr>(Attr) || isa<CUDADeviceAttr>(Attr) || + isa<CUDAGlobalAttr>(Attr))) { + // CUDA target attributes are part of function signature for + // overloading purposes and must not be merged. + return false; + } else if (const auto *MA = dyn_cast<MinSizeAttr>(Attr)) NewAttr = S.mergeMinSizeAttr(D, MA->getRange(), AttrSpellingListIndex); else if (const auto *OA = dyn_cast<OptimizeNoneAttr>(Attr)) NewAttr = S.mergeOptimizeNoneAttr(D, OA->getRange(), AttrSpellingListIndex); Modified: cfe/trunk/test/SemaCUDA/function-overload.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=281406&r1=281405&r2=281406&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/function-overload.cu (original) +++ cfe/trunk/test/SemaCUDA/function-overload.cu Tue Sep 13 17:16:30 2016 @@ -379,3 +379,14 @@ __host__ __device__ void test_host_devic HostReturnTy ret3 = host_only_function(1); HostReturnTy2 ret4 = host_only_function(1.0f); } + +// Verify that we allow overloading function templates. +template <typename T> __host__ T template_overload(const T &a) { return a; }; +template <typename T> __device__ T template_overload(const T &a) { return a; }; + +__host__ void test_host_template_overload() { + template_overload(1); // OK. Attribute-based overloading picks __host__ variant. +} +__device__ void test_device_template_overload() { + template_overload(1); // OK. Attribute-based overloading picks __device__ variant. +} Added: cfe/trunk/test/SemaCUDA/target_attr_inheritance.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/target_attr_inheritance.cu?rev=281406&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/target_attr_inheritance.cu (added) +++ cfe/trunk/test/SemaCUDA/target_attr_inheritance.cu Tue Sep 13 17:16:30 2016 @@ -0,0 +1,29 @@ +// Verifies correct inheritance of target attributes during template +// instantiation and specialization. + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s + +#include "Inputs/cuda.h" + +// Function must inherit target attributes during instantiation, but not during +// specialization. +template <typename T> __host__ __device__ T function_template(const T &a); + +// Specialized functions have their own attributes. +// expected-note@+1 {{candidate function not viable: call to __host__ function from __device__ function}} +template <> __host__ float function_template<float>(const float &from); + +// expected-note@+1 {{candidate function not viable: call to __device__ function from __host__ function}} +template <> __device__ double function_template<double>(const double &from); + +__host__ void hf() { + function_template<float>(1.0f); // OK. Specialization is __host__. + function_template<double>(2.0); // expected-error {{no matching function for call to 'function_template'}} + function_template(1); // OK. Instantiated function template is HD. +} +__device__ void df() { + function_template<float>(3.0f); // expected-error {{no matching function for call to 'function_template'}} + function_template<double>(4.0); // OK. Specialization is __device__. + function_template(1); // OK. Instantiated function template is HD. +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits