r354909 - [AMDGPU] Allow using integral non-type template parameters
Michael Liao via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 26 10:49:36 PST 2019
Author: hliao
Date: Tue Feb 26 10:49:36 2019
New Revision: 354909
URL: http://llvm.org/viewvc/llvm-project?rev=354909&view=rev
Log:
[AMDGPU] Allow using integral non-type template parameters
Summary:
- Allow using integral non-type template parameters in the following
attributes
__attribute__((amdgpu_flat_work_group_size(<min>, <max>)))
__attribute__((amdgpu_waves_per_eu(<min>[, <max>])))
Reviewers: kzhuravl, yaxunl
Subscribers: jvesely, wdng, nhaehnle, dstuttard, tpr, t-tye, jdoerfert, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D58623
Modified:
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/CodeGen/TargetInfo.cpp
cfe/trunk/lib/Sema/SemaDeclAttr.cpp
cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp
cfe/trunk/test/SemaCUDA/amdgpu-attrs.cu
cfe/trunk/test/SemaOpenCL/amdgpu-attrs.cl
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Tue Feb 26 10:49:36 2019
@@ -1484,14 +1484,14 @@ def RISCVInterrupt : InheritableAttr, Ta
def AMDGPUFlatWorkGroupSize : InheritableAttr {
let Spellings = [Clang<"amdgpu_flat_work_group_size", 0>];
- let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max">];
+ let Args = [ExprArgument<"Min">, ExprArgument<"Max">];
let Documentation = [AMDGPUFlatWorkGroupSizeDocs];
let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
}
def AMDGPUWavesPerEU : InheritableAttr {
let Spellings = [Clang<"amdgpu_waves_per_eu", 0>];
- let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max", 1>];
+ let Args = [ExprArgument<"Min">, ExprArgument<"Max", 1>];
let Documentation = [AMDGPUWavesPerEUDocs];
let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
}
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Tue Feb 26 10:49:36 2019
@@ -8674,6 +8674,16 @@ public:
void AddXConsumedAttr(Decl *D, SourceRange SR, unsigned SpellingIndex,
RetainOwnershipKind K, bool IsTemplateInstantiation);
+ /// addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size
+ /// attribute to a particular declaration.
+ void addAMDGPUFlatWorkGroupSizeAttr(SourceRange AttrRange, Decl *D, Expr *Min,
+ Expr *Max, unsigned SpellingListIndex);
+
+ /// addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a
+ /// particular declaration.
+ void addAMDGPUWavesPerEUAttr(SourceRange AttrRange, Decl *D, Expr *Min,
+ Expr *Max, unsigned SpellingListIndex);
+
bool checkNSReturnsRetainedReturnType(SourceLocation loc, QualType type);
//===--------------------------------------------------------------------===//
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Tue Feb 26 10:49:36 2019
@@ -7797,8 +7797,16 @@ void AMDGPUTargetCodeGenInfo::setTargetA
const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
if (ReqdWGS || FlatWGS) {
- unsigned Min = FlatWGS ? FlatWGS->getMin() : 0;
- unsigned Max = FlatWGS ? FlatWGS->getMax() : 0;
+ unsigned Min = 0;
+ unsigned Max = 0;
+ if (FlatWGS) {
+ Min = FlatWGS->getMin()
+ ->EvaluateKnownConstInt(M.getContext())
+ .getExtValue();
+ Max = FlatWGS->getMax()
+ ->EvaluateKnownConstInt(M.getContext())
+ .getExtValue();
+ }
if (ReqdWGS && Min == 0 && Max == 0)
Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim();
@@ -7812,8 +7820,12 @@ void AMDGPUTargetCodeGenInfo::setTargetA
}
if (const auto *Attr = FD->getAttr<AMDGPUWavesPerEUAttr>()) {
- unsigned Min = Attr->getMin();
- unsigned Max = Attr->getMax();
+ unsigned Min =
+ Attr->getMin()->EvaluateKnownConstInt(M.getContext()).getExtValue();
+ unsigned Max = Attr->getMax() ? Attr->getMax()
+ ->EvaluateKnownConstInt(M.getContext())
+ .getExtValue()
+ : 0;
if (Min != 0) {
assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max");
Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Tue Feb 26 10:49:36 2019
@@ -245,11 +245,11 @@ static bool checkUInt32Argument(Sema &S,
!Expr->isIntegerConstantExpr(I, S.Context)) {
if (Idx != UINT_MAX)
S.Diag(getAttrLoc(AI), diag::err_attribute_argument_n_type)
- << AI << Idx << AANT_ArgumentIntegerConstant
+ << &AI << Idx << AANT_ArgumentIntegerConstant
<< Expr->getSourceRange();
else
S.Diag(getAttrLoc(AI), diag::err_attribute_argument_type)
- << AI << AANT_ArgumentIntegerConstant << Expr->getSourceRange();
+ << &AI << AANT_ArgumentIntegerConstant << Expr->getSourceRange();
return false;
}
@@ -261,7 +261,7 @@ static bool checkUInt32Argument(Sema &S,
if (StrictlyUnsigned && I.isSigned() && I.isNegative()) {
S.Diag(getAttrLoc(AI), diag::err_attribute_requires_positive_integer)
- << AI << /*non-negative*/ 1;
+ << &AI << /*non-negative*/ 1;
return false;
}
@@ -5853,57 +5853,115 @@ static void handleInterruptAttr(Sema &S,
}
}
-static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D,
- const ParsedAttr &AL) {
+static bool
+checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
+ const AMDGPUFlatWorkGroupSizeAttr &Attr) {
+ // Accept template arguments for now as they depend on something else.
+ // We'll get to check them when they eventually get instantiated.
+ if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
+ return false;
+
uint32_t Min = 0;
- Expr *MinExpr = AL.getArgAsExpr(0);
- if (!checkUInt32Argument(S, AL, MinExpr, Min))
- return;
+ if (!checkUInt32Argument(S, Attr, MinExpr, Min, 0))
+ return true;
uint32_t Max = 0;
- Expr *MaxExpr = AL.getArgAsExpr(1);
- if (!checkUInt32Argument(S, AL, MaxExpr, Max))
- return;
+ if (!checkUInt32Argument(S, Attr, MaxExpr, Max, 1))
+ return true;
if (Min == 0 && Max != 0) {
- S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 0;
- return;
+ S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+ << &Attr << 0;
+ return true;
}
if (Min > Max) {
- S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
- return;
+ S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+ << &Attr << 1;
+ return true;
}
- D->addAttr(::new (S.Context)
- AMDGPUFlatWorkGroupSizeAttr(AL.getLoc(), S.Context, Min, Max,
- AL.getAttributeSpellingListIndex()));
+ return false;
}
-static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
- uint32_t Min = 0;
- Expr *MinExpr = AL.getArgAsExpr(0);
- if (!checkUInt32Argument(S, AL, MinExpr, Min))
+void Sema::addAMDGPUFlatWorkGroupSizeAttr(SourceRange AttrRange, Decl *D,
+ Expr *MinExpr, Expr *MaxExpr,
+ unsigned SpellingListIndex) {
+ AMDGPUFlatWorkGroupSizeAttr TmpAttr(AttrRange, Context, MinExpr, MaxExpr,
+ SpellingListIndex);
+
+ if (checkAMDGPUFlatWorkGroupSizeArguments(*this, MinExpr, MaxExpr, TmpAttr))
return;
+ D->addAttr(::new (Context) AMDGPUFlatWorkGroupSizeAttr(
+ AttrRange, Context, MinExpr, MaxExpr, SpellingListIndex));
+}
+
+static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D,
+ const ParsedAttr &AL) {
+ Expr *MinExpr = AL.getArgAsExpr(0);
+ Expr *MaxExpr = AL.getArgAsExpr(1);
+
+ S.addAMDGPUFlatWorkGroupSizeAttr(AL.getRange(), D, MinExpr, MaxExpr,
+ AL.getAttributeSpellingListIndex());
+}
+
+static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
+ Expr *MaxExpr,
+ const AMDGPUWavesPerEUAttr &Attr) {
+ if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
+ (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
+ return true;
+
+ // Accept template arguments for now as they depend on something else.
+ // We'll get to check them when they eventually get instantiated.
+ if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
+ return false;
+
+ uint32_t Min = 0;
+ if (!checkUInt32Argument(S, Attr, MinExpr, Min, 0))
+ return true;
+
uint32_t Max = 0;
- if (AL.getNumArgs() == 2) {
- Expr *MaxExpr = AL.getArgAsExpr(1);
- if (!checkUInt32Argument(S, AL, MaxExpr, Max))
- return;
- }
+ if (MaxExpr && !checkUInt32Argument(S, Attr, MaxExpr, Max, 1))
+ return true;
if (Min == 0 && Max != 0) {
- S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 0;
- return;
+ S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+ << &Attr << 0;
+ return true;
}
if (Max != 0 && Min > Max) {
- S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
- return;
+ S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+ << &Attr << 1;
+ return true;
}
- D->addAttr(::new (S.Context)
- AMDGPUWavesPerEUAttr(AL.getLoc(), S.Context, Min, Max,
- AL.getAttributeSpellingListIndex()));
+ return false;
+}
+
+void Sema::addAMDGPUWavesPerEUAttr(SourceRange AttrRange, Decl *D,
+ Expr *MinExpr, Expr *MaxExpr,
+ unsigned SpellingListIndex) {
+ AMDGPUWavesPerEUAttr TmpAttr(AttrRange, Context, MinExpr, MaxExpr,
+ SpellingListIndex);
+
+ if (checkAMDGPUWavesPerEUArguments(*this, MinExpr, MaxExpr, TmpAttr))
+ return;
+
+ D->addAttr(::new (Context) AMDGPUWavesPerEUAttr(AttrRange, Context, MinExpr,
+ MaxExpr, SpellingListIndex));
+}
+
+static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+ if (!checkAttributeAtLeastNumArgs(S, AL, 1) ||
+ !checkAttributeAtMostNumArgs(S, AL, 2))
+ return;
+
+ Expr *MinExpr = AL.getArgAsExpr(0);
+ Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
+
+ S.addAMDGPUWavesPerEUAttr(AL.getRange(), D, MinExpr, MaxExpr,
+ AL.getAttributeSpellingListIndex());
}
static void handleAMDGPUNumSGPRAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
Modified: cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp Tue Feb 26 10:49:36 2019
@@ -344,6 +344,51 @@ static void instantiateOMPDeclareSimdDec
Attr.getRange());
}
+static void instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
+ Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+ const AMDGPUFlatWorkGroupSizeAttr &Attr, Decl *New) {
+ // Both min and max expression are constant expressions.
+ EnterExpressionEvaluationContext Unevaluated(
+ S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+
+ ExprResult Result = S.SubstExpr(Attr.getMin(), TemplateArgs);
+ if (Result.isInvalid())
+ return;
+ Expr *MinExpr = Result.getAs<Expr>();
+
+ Result = S.SubstExpr(Attr.getMax(), TemplateArgs);
+ if (Result.isInvalid())
+ return;
+ Expr *MaxExpr = Result.getAs<Expr>();
+
+ S.addAMDGPUFlatWorkGroupSizeAttr(Attr.getLocation(), New, MinExpr, MaxExpr,
+ Attr.getSpellingListIndex());
+}
+
+static void instantiateDependentAMDGPUWavesPerEUAttr(
+ Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+ const AMDGPUWavesPerEUAttr &Attr, Decl *New) {
+ // Both min and max expression are constant expressions.
+ EnterExpressionEvaluationContext Unevaluated(
+ S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+
+ ExprResult Result = S.SubstExpr(Attr.getMin(), TemplateArgs);
+ if (Result.isInvalid())
+ return;
+ Expr *MinExpr = Result.getAs<Expr>();
+
+ Expr *MaxExpr = nullptr;
+ if (auto Max = Attr.getMax()) {
+ Result = S.SubstExpr(Max, TemplateArgs);
+ if (Result.isInvalid())
+ return;
+ MaxExpr = Result.getAs<Expr>();
+ }
+
+ S.addAMDGPUWavesPerEUAttr(Attr.getLocation(), New, MinExpr, MaxExpr,
+ Attr.getSpellingListIndex());
+}
+
void Sema::InstantiateAttrsForDecl(
const MultiLevelTemplateArgumentList &TemplateArgs, const Decl *Tmpl,
Decl *New, LateInstantiatedAttrVec *LateAttrs,
@@ -437,6 +482,18 @@ void Sema::InstantiateAttrs(const MultiL
continue;
}
+ if (const AMDGPUFlatWorkGroupSizeAttr *AMDGPUFlatWorkGroupSize =
+ dyn_cast<AMDGPUFlatWorkGroupSizeAttr>(TmplAttr)) {
+ instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
+ *this, TemplateArgs, *AMDGPUFlatWorkGroupSize, New);
+ }
+
+ if (const AMDGPUWavesPerEUAttr *AMDGPUFlatWorkGroupSize =
+ dyn_cast<AMDGPUWavesPerEUAttr>(TmplAttr)) {
+ instantiateDependentAMDGPUWavesPerEUAttr(*this, TemplateArgs,
+ *AMDGPUFlatWorkGroupSize, New);
+ }
+
// Existing DLL attribute on the instantiation takes precedence.
if (TmplAttr->getKind() == attr::DLLExport ||
TmplAttr->getKind() == attr::DLLImport) {
Modified: cfe/trunk/test/SemaCUDA/amdgpu-attrs.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/amdgpu-attrs.cu?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/amdgpu-attrs.cu (original)
+++ cfe/trunk/test/SemaCUDA/amdgpu-attrs.cu Tue Feb 26 10:49:36 2019
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
#include "Inputs/cuda.h"
@@ -78,3 +78,119 @@ __global__ void vec_type_hint_int() {}
// expected-error at +2{{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel function}}
__attribute__((intel_reqd_sub_group_size(64)))
__global__ void intel_reqd_sub_group_size_64() {}
+
+// expected-error at +1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("32", 64)))
+__global__ void non_int_min_flat_work_group_size_32_64() {}
+// expected-error at +1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, "64")))
+__global__ void non_int_max_flat_work_group_size_32_64() {}
+
+int nc_min = 32, nc_max = 64;
+// expected-error at +1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(nc_min, 64)))
+__global__ void non_cint_min_flat_work_group_size_32_64() {}
+// expected-error at +1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, nc_max)))
+__global__ void non_cint_max_flat_work_group_size_32_64() {}
+
+const int c_min = 16, c_max = 32;
+__attribute__((amdgpu_flat_work_group_size(c_min * 2, 64)))
+__global__ void cint_min_flat_work_group_size_32_64() {}
+__attribute__((amdgpu_flat_work_group_size(32, c_max * 2)))
+__global__ void cint_max_flat_work_group_size_32_64() {}
+
+// expected-error at +3{{'T' does not refer to a value}}
+// expected-note at +1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_flat_work_group_size(T, 64)))
+__global__ void template_class_min_flat_work_group_size_32_64() {}
+// expected-error at +3{{'T' does not refer to a value}}
+// expected-note at +1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_flat_work_group_size(32, T)))
+__global__ void template_class_max_flat_work_group_size_32_64() {}
+
+template<unsigned a, unsigned b>
+__attribute__((amdgpu_flat_work_group_size(a, b)))
+__global__ void template_flat_work_group_size_32_64() {}
+template __global__ void template_flat_work_group_size_32_64<32, 64>();
+
+template<unsigned a, unsigned b, unsigned c>
+__attribute__((amdgpu_flat_work_group_size(a + b, b + c)))
+__global__ void template_complex_flat_work_group_size_32_64() {}
+template __global__ void template_complex_flat_work_group_size_32_64<16, 16, 48>();
+
+unsigned ipow2(unsigned n) { return n == 0 ? 1 : 2 * ipow2(n - 1); }
+constexpr unsigned ce_ipow2(unsigned n) { return n == 0 ? 1 : 2 * ce_ipow2(n - 1); }
+
+__attribute__((amdgpu_flat_work_group_size(ce_ipow2(5), ce_ipow2(6))))
+__global__ void cexpr_flat_work_group_size_32_64() {}
+// expected-error at +1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(ipow2(5), 64)))
+__global__ void non_cexpr_min_flat_work_group_size_32_64() {}
+// expected-error at +1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, ipow2(6))))
+__global__ void non_cexpr_max_flat_work_group_size_32_64() {}
+
+// expected-error at +1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu("2")))
+__global__ void non_int_min_waves_per_eu_2() {}
+// expected-error at +1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, "4")))
+__global__ void non_int_max_waves_per_eu_2_4() {}
+
+// expected-error at +1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(nc_min)))
+__global__ void non_cint_min_waves_per_eu_2() {}
+// expected-error at +1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, nc_max)))
+__global__ void non_cint_min_waves_per_eu_2_4() {}
+
+__attribute__((amdgpu_waves_per_eu(c_min / 8)))
+__global__ void cint_min_waves_per_eu_2() {}
+__attribute__((amdgpu_waves_per_eu(c_min / 8, c_max / 8)))
+__global__ void cint_min_waves_per_eu_2_4() {}
+
+// expected-error at +3{{'T' does not refer to a value}}
+// expected-note at +1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_waves_per_eu(T)))
+__global__ void cint_min_waves_per_eu_2() {}
+// expected-error at +3{{'T' does not refer to a value}}
+// expected-note at +1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_waves_per_eu(2, T)))
+__global__ void cint_min_waves_per_eu_2_4() {}
+
+template<unsigned a>
+__attribute__((amdgpu_waves_per_eu(a)))
+__global__ void template_waves_per_eu_2() {}
+template __global__ void template_waves_per_eu_2<2>();
+
+template<unsigned a, unsigned b>
+__attribute__((amdgpu_waves_per_eu(a, b)))
+__global__ void template_waves_per_eu_2_4() {}
+template __global__ void template_waves_per_eu_2_4<2, 4>();
+
+template<unsigned a, unsigned b, unsigned c>
+__attribute__((amdgpu_waves_per_eu(a + b, c - b)))
+__global__ void template_complex_waves_per_eu_2_4() {}
+template __global__ void template_complex_waves_per_eu_2_4<1, 1, 5>();
+
+// expected-error at +2{{expression contains unexpanded parameter pack 'Args'}}
+template<unsigned... Args>
+__attribute__((amdgpu_waves_per_eu(Args)))
+__global__ void template_waves_per_eu_2() {}
+template __global__ void template_waves_per_eu_2<2, 4>();
+
+__attribute__((amdgpu_waves_per_eu(ce_ipow2(1))))
+__global__ void cexpr_waves_per_eu_2() {}
+__attribute__((amdgpu_waves_per_eu(ce_ipow2(1), ce_ipow2(2))))
+__global__ void cexpr_waves_per_eu_2_4() {}
+// expected-error at +1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(ipow2(1))))
+__global__ void non_cexpr_waves_per_eu_2() {}
+// expected-error at +1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, ipow2(2))))
+__global__ void non_cexpr_waves_per_eu_2_4() {}
Modified: cfe/trunk/test/SemaOpenCL/amdgpu-attrs.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/amdgpu-attrs.cl?rev=354909&r1=354908&r2=354909&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/amdgpu-attrs.cl (original)
+++ cfe/trunk/test/SemaOpenCL/amdgpu-attrs.cl Tue Feb 26 10:49:36 2019
@@ -27,12 +27,12 @@ __attribute__((amdgpu_waves_per_eu(2, 4)
__attribute__((amdgpu_num_sgpr(32))) void func_num_sgpr_32() {} // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
__attribute__((amdgpu_num_vgpr(64))) void func_num_vgpr_64() {} // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
-__attribute__((amdgpu_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
-__attribute__((amdgpu_flat_work_group_size(32, "ABC"))) kernel void kernel_flat_work_group_size_32_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
-__attribute__((amdgpu_flat_work_group_size("ABC", 64))) kernel void kernel_flat_work_group_size_ABC_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
-__attribute__((amdgpu_waves_per_eu("ABC"))) kernel void kernel_waves_per_eu_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
-__attribute__((amdgpu_waves_per_eu(2, "ABC"))) kernel void kernel_waves_per_eu_2_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
-__attribute__((amdgpu_waves_per_eu("ABC", 4))) kernel void kernel_waves_per_eu_ABC_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, "ABC"))) kernel void kernel_flat_work_group_size_32_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("ABC", 64))) kernel void kernel_flat_work_group_size_ABC_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu("ABC"))) kernel void kernel_waves_per_eu_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, "ABC"))) kernel void kernel_waves_per_eu_2_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu("ABC", 4))) kernel void kernel_waves_per_eu_ABC_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
__attribute__((amdgpu_num_sgpr("ABC"))) kernel void kernel_num_sgpr_ABC() {} // expected-error {{'amdgpu_num_sgpr' attribute requires an integer constant}}
__attribute__((amdgpu_num_vgpr("ABC"))) kernel void kernel_num_vgpr_ABC() {} // expected-error {{'amdgpu_num_vgpr' attribute requires an integer constant}}
More information about the cfe-commits
mailing list