Skip to content

Commit 5b37b04

Browse files
AlexMacleanAlexMaclean
committed
address comments, fixup test
1 parent b637f2a commit 5b37b04

File tree

4 files changed

+29
-24
lines changed

4 files changed

+29
-24
lines changed

‎clang/test/OpenMP/ompx_attributes_codegen.cpp

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -11,13 +11,13 @@
1111

1212
// Check that the target attributes are set on the generated kernel
1313
void func() {
14-
// AMD: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l18(ptr {{[^,]+}}) #0
15-
// AMD: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l20(ptr {{[^,]+}})
16-
// AMD: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l22(ptr {{[^,]+}}) #4
14+
// AMD: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l22(ptr {{[^,]+}}) #0
15+
// AMD: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l24(ptr {{[^,]+}})
16+
// AMD: amdgpu_kernel void @__omp_offloading[[HASH:.*]]_l26(ptr {{[^,]+}}) #4
1717

18-
// NVIDIA: ptx_kernel void @__omp_offloading[[HASH:.*]]_l18(ptr {{[^,]+}}) #[[ATTR0:[0-9]+]]
19-
// NVIDIA: ptx_kernel void @__omp_offloading[[HASH:.*]]_l20(ptr {{[^,]+}}) #[[ATTR1:[0-9]+]]
20-
// NVIDIA: ptx_kernel void @__omp_offloading[[HASH:.*]]_l22(ptr {{[^,]+}}) #[[ATTR2:[0-9]+]]
18+
// NVIDIA: ptx_kernel void @__omp_offloading[[HASH:.*]]_l22(ptr {{[^,]+}}) #[[ATTR0:[0-9]+]]
19+
// NVIDIA: ptx_kernel void @__omp_offloading[[HASH:.*]]_l24(ptr {{[^,]+}}) #[[ATTR1:[0-9]+]]
20+
// NVIDIA: ptx_kernel void @__omp_offloading[[HASH:.*]]_l26(ptr {{[^,]+}}) #[[ATTR2:[0-9]+]]
2121

2222
#pragma omp target ompx_attribute([[clang::amdgpu_flat_work_group_size(10, 20)]])
2323
{}
@@ -39,11 +39,11 @@ void func() {
3939

4040
// It is unclear if we should use the AMD annotations for other targets, we do for now.
4141
// NVIDIA: attributes #[[ATTR0]]
42-
// NVIDIA-SAME: "omp_target_thread_limit"="20"
4342
// NVIDIA-SAME: "nvvm.maxntid"="20"
43+
// NVIDIA-SAME: "omp_target_thread_limit"="20"
4444
// NVIDIA: attributes #[[ATTR1]]
45-
// NVIDIA-SAME: "omp_target_thread_limit"="45"
4645
// NVIDIA-SAME: "nvvm.maxntid"="45"
46+
// NVIDIA-SAME: "omp_target_thread_limit"="45"
4747
// NVIDIA: attributes #[[ATTR2]]
48-
// NVIDIA-SAME: "omp_target_thread_limit"="17"
4948
// NVIDIA-SAME: "nvvm.maxntid"="17"
49+
// NVIDIA-SAME: "omp_target_thread_limit"="17"

‎llvm/lib/IR/AutoUpgrade.cpp

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -5033,6 +5033,8 @@ static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC,
50335033
unsigned Length = 0;
50345034

50355035
if (F->hasFnAttribute(Attr)) {
5036+
// We expect the existing attribute to have the form "x[,y[,z]]". Here we
5037+
// parse these elements placing them into Vect3
50365038
StringRef S = F->getFnAttribute(Attr).getValueAsString();
50375039
for (; Length < 3 && !S.empty(); Length++) {
50385040
auto [Part, Rest] = S.split(',');
@@ -5041,19 +5043,22 @@ static void upgradeNVVMFnVectorAttr(const StringRef Attr, const char DimC,
50415043
}
50425044
}
50435045

5044-
const uint64_t VInt = mdconst::extract<ConstantInt>(V)->getZExtValue();
5045-
const std::string VStr = llvm::utostr(VInt);
5046-
50475046
const unsigned Dim = DimC - 'x';
50485047
assert(Dim >= 0 && Dim < 3 && "Unexpected dim char");
50495048

5049+
const uint64_t VInt = mdconst::extract<ConstantInt>(V)->getZExtValue();
5050+
const std::string VStr = llvm::utostr(VInt);
50505051
Vect3[Dim] = VStr;
50515052
Length = std::max(Length, Dim + 1);
50525053

50535054
const std::string NewAttr = llvm::join(ArrayRef(Vect3, Length), ",");
50545055
F->addFnAttr(Attr, NewAttr);
50555056
}
50565057

5058+
static inline bool isXYZ(StringRef S) {
5059+
return S == "x" || S == "y" || S == "z";
5060+
}
5061+
50575062
bool static upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K,
50585063
const Metadata *V) {
50595064
if (K == "kernel") {
@@ -5092,15 +5097,15 @@ bool static upgradeSingleNVVMAnnotation(GlobalValue *GV, StringRef K,
50925097
cast<Function>(GV)->addFnAttr("nvvm.maxnreg", llvm::utostr(CV));
50935098
return true;
50945099
}
5095-
if (K.consume_front("maxntid") && (K == "x" || K == "y" || K == "z")) {
5100+
if (K.consume_front("maxntid") && isXYZ(K)) {
50965101
upgradeNVVMFnVectorAttr("nvvm.maxntid", K[0], GV, V);
50975102
return true;
50985103
}
5099-
if (K.consume_front("reqntid") && (K == "x" || K == "y" || K == "z")) {
5104+
if (K.consume_front("reqntid") && isXYZ(K)) {
51005105
upgradeNVVMFnVectorAttr("nvvm.reqntid", K[0], GV, V);
51015106
return true;
51025107
}
5103-
if (K.consume_front("cluster_dim_") && (K == "x" || K == "y" || K == "z")) {
5108+
if (K.consume_front("cluster_dim_") && isXYZ(K)) {
51045109
upgradeNVVMFnVectorAttr("nvvm.cluster_dim", K[0], GV, V);
51055110
return true;
51065111
}

‎llvm/lib/Target/NVPTX/NVPTXUtilities.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,7 @@
2222
#include "llvm/IR/Module.h"
2323
#include "llvm/Support/Alignment.h"
2424
#include "llvm/Support/Mutex.h"
25+
#include <cstdint>
2526
#include <cstring>
2627
#include <map>
2728
#include <mutex>
@@ -204,6 +205,8 @@ static SmallVector<unsigned, 3> getFnAttrParsedVector(const Function &F,
204205
auto &Ctx = F.getContext();
205206

206207
if (F.hasFnAttribute(Attr)) {
208+
// We expect the attribute value to be of the form "x[,y[,z]]", where x, y,
209+
// and z are unsigned values.
207210
StringRef S = F.getFnAttribute(Attr).getValueAsString();
208211
for (unsigned I = 0; I < 3 && !S.empty(); I++) {
209212
auto [First, Rest] = S.split(",");
@@ -218,14 +221,11 @@ static SmallVector<unsigned, 3> getFnAttrParsedVector(const Function &F,
218221
return V;
219222
}
220223

221-
static std::optional<unsigned> getVectorProduct(ArrayRef<unsigned> V) {
224+
static std::optional<uint64_t> getVectorProduct(ArrayRef<unsigned> V) {
222225
if (V.empty())
223226
return std::nullopt;
224227

225-
unsigned Product = 1;
226-
for (const unsigned E : V)
227-
Product *= E;
228-
return Product;
228+
return std::accumulate(V.begin(), V.end(), 1, std::multiplies<uint64_t>{});
229229
}
230230

231231
bool isParamGridConstant(const Value &V) {
@@ -298,7 +298,7 @@ SmallVector<unsigned, 3> getClusterDim(const Function &F) {
298298
return getFnAttrParsedVector(F, "nvvm.cluster_dim");
299299
}
300300

301-
std::optional<unsigned> getOverallMaxNTID(const Function &F) {
301+
std::optional<uint64_t> getOverallMaxNTID(const Function &F) {
302302
// Note: The semantics here are a bit strange. The PTX ISA states the
303303
// following (11.4.2. Performance-Tuning Directives: .maxntid):
304304
//
@@ -309,7 +309,7 @@ std::optional<unsigned> getOverallMaxNTID(const Function &F) {
309309
return getVectorProduct(MaxNTID);
310310
}
311311

312-
std::optional<unsigned> getOverallReqNTID(const Function &F) {
312+
std::optional<uint64_t> getOverallReqNTID(const Function &F) {
313313
// Note: The semantics here are a bit strange. See getMaxNTID.
314314
const auto ReqNTID = getReqNTID(F);
315315
return getVectorProduct(ReqNTID);

‎llvm/lib/Target/NVPTX/NVPTXUtilities.h

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -52,8 +52,8 @@ SmallVector<unsigned, 3> getMaxNTID(const Function &);
5252
SmallVector<unsigned, 3> getReqNTID(const Function &);
5353
SmallVector<unsigned, 3> getClusterDim(const Function &);
5454

55-
std::optional<unsigned> getOverallMaxNTID(const Function &);
56-
std::optional<unsigned> getOverallReqNTID(const Function &);
55+
std::optional<uint64_t> getOverallMaxNTID(const Function &);
56+
std::optional<uint64_t> getOverallReqNTID(const Function &);
5757

5858
std::optional<unsigned> getMaxClusterRank(const Function &);
5959
std::optional<unsigned> getMinCTASm(const Function &);

0 commit comments

Comments
 (0)