Skip to content

Commit 716c8f0

Browse files
committed
removing const
1 parent 1ac6c2b commit 716c8f0

File tree

2 files changed

+33
-35
lines changed

2 files changed

+33
-35
lines changed

clang/test/CodeGen/builtins-nvptx.c

Lines changed: 11 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1,37 +1,37 @@
11
// REQUIRES: nvptx-registered-target
22
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_70 -target-feature +ptx63 \
3-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
3+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
44
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX63_SM70 -check-prefix=LP64 %s
55
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_80 -target-feature +ptx70 \
6-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
6+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
77
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 -check-prefix=LP32 %s
88
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_80 -target-feature +ptx70 \
9-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
9+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
1010
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 -check-prefix=LP64 %s
1111
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_60 -target-feature +ptx62 \
12-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
12+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
1313
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP32 %s
1414
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_60 -target-feature +ptx62 \
15-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
15+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
1616
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s
1717
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_61 -target-feature +ptx62 \
18-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
18+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
1919
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s
2020
// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_53 -target-feature +ptx62 \
2121
// RUN: -DERROR_CHECK -fcuda-is-device -S -o /dev/null -x cuda -verify %s
2222
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_86 -target-feature +ptx72 \
23-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
23+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
2424
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 -check-prefix=LP32 %s
2525
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_86 -target-feature +ptx72 \
26-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
26+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
2727
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 -check-prefix=LP64 %s
2828
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_89 -target-feature +ptx81 \
29-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
29+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
3030
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s
3131
// ### The last run to check with the highest SM and PTX version available
3232
// ### to make sure target builtins are still accepted.
3333
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_100a -target-feature +ptx87 \
34-
// RUN: -fcuda-is-device -emit-llvm -o - -x cuda %s \
34+
// RUN: -disable-llvm-optzns -fcuda-is-device -emit-llvm -o - -x cuda %s \
3535
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX81_SM89 %s
3636

3737
#define __device__ __attribute__((device))
@@ -61,6 +61,7 @@ __device__ bool reflect() {
6161

6262
unsigned x = __nvvm_reflect("__CUDA_ARCH");
6363
return x >= 700;
64+
6465
}
6566

6667
__device__ int read_ntid() {

llvm/lib/Target/NVPTX/NVVMReflect.cpp

Lines changed: 22 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -61,14 +61,14 @@ class NVVMReflect {
6161
// Map from reflect function call arguments to the value to replace the call
6262
// with. Should include __CUDA_FTZ and __CUDA_ARCH values.
6363
StringMap<unsigned> ReflectMap;
64-
bool handleReflectFunction(Module &M, const StringRef ReflectName);
64+
bool handleReflectFunction(Module &M, StringRef ReflectName);
6565
void populateReflectMap(Module &M);
66-
void foldReflectCall(CallInst *const Call, Constant *const NewValue);
66+
void foldReflectCall(CallInst *Call, Constant *NewValue);
6767

6868
public:
6969
// __CUDA_FTZ is assigned in `runOnModule` by checking nvvm-reflect-ftz module
7070
// metadata.
71-
explicit NVVMReflect(const unsigned SmVersion)
71+
explicit NVVMReflect(unsigned SmVersion)
7272
: ReflectMap({{CUDA_ARCH_NAME, SmVersion * 10}}) {}
7373
bool runOnModule(Module &M);
7474
};
@@ -78,13 +78,13 @@ class NVVMReflectLegacyPass : public ModulePass {
7878

7979
public:
8080
static char ID;
81-
NVVMReflectLegacyPass(const unsigned SmVersion)
81+
NVVMReflectLegacyPass(unsigned SmVersion)
8282
: ModulePass(ID), Impl(SmVersion) {}
8383
bool runOnModule(Module &M) override;
8484
};
8585
} // namespace
8686

87-
ModulePass *llvm::createNVVMReflectPass(const unsigned SmVersion) {
87+
ModulePass *llvm::createNVVMReflectPass(unsigned SmVersion) {
8888
return new NVVMReflectLegacyPass(SmVersion);
8989
}
9090

@@ -109,11 +109,11 @@ static cl::list<std::string> ReflectList(
109109
// Set the VarMap with, first, the value of __CUDA_FTZ from module metadata, and
110110
// then the key/value pairs from the command line.
111111
void NVVMReflect::populateReflectMap(Module &M) {
112-
if (const auto *const Flag = mdconst::extract_or_null<ConstantInt>(
112+
if (auto *Flag = mdconst::extract_or_null<ConstantInt>(
113113
M.getModuleFlag(CUDA_FTZ_MODULE_NAME)))
114114
ReflectMap[CUDA_FTZ_NAME] = Flag->getSExtValue();
115115

116-
for (const auto &Option : ReflectList) {
116+
for (auto &Option : ReflectList) {
117117
LLVM_DEBUG(dbgs() << "ReflectOption : " << Option << "\n");
118118
StringRef OptionRef(Option);
119119
auto [Name, Val] = OptionRef.split('=');
@@ -135,43 +135,42 @@ void NVVMReflect::populateReflectMap(Module &M) {
135135
/// Process a reflect function by finding all its calls and replacing them with
136136
/// appropriate constant values. For __CUDA_FTZ, uses the module flag value.
137137
/// For __CUDA_ARCH, uses SmVersion * 10. For all other strings, uses 0.
138-
bool NVVMReflect::handleReflectFunction(Module &M,
139-
const StringRef ReflectName) {
140-
Function *const F = M.getFunction(ReflectName);
138+
bool NVVMReflect::handleReflectFunction(Module &M, StringRef ReflectName) {
139+
Function *F = M.getFunction(ReflectName);
141140
if (!F)
142141
return false;
143142
assert(F->isDeclaration() && "_reflect function should not have a body");
144143
assert(F->getReturnType()->isIntegerTy() &&
145144
"_reflect's return type should be integer");
146145

147146
const bool Changed = F->getNumUses() > 0;
148-
for (User *const U : make_early_inc_range(F->users())) {
147+
for (User *U : make_early_inc_range(F->users())) {
149148
// Reflect function calls look like:
150149
// @arch = private unnamed_addr addrspace(1) constant [12 x i8]
151150
// c"__CUDA_ARCH\00" call i32 @__nvvm_reflect(ptr addrspacecast (ptr
152151
// addrspace(1) @arch to ptr)) We need to extract the string argument from
153152
// the call (i.e. "__CUDA_ARCH")
154-
auto *const Call = dyn_cast<CallInst>(U);
153+
auto *Call = dyn_cast<CallInst>(U);
155154
if (!Call)
156155
report_fatal_error(
157156
"__nvvm_reflect can only be used in a call instruction");
158157
if (Call->getNumOperands() != 2)
159158
report_fatal_error("__nvvm_reflect requires exactly one argument");
160159

161-
const auto *const GlobalStr =
160+
auto *GlobalStr =
162161
dyn_cast<Constant>(Call->getArgOperand(0)->stripPointerCasts());
163162
if (!GlobalStr)
164163
report_fatal_error("__nvvm_reflect argument must be a constant string");
165164

166-
const auto *const ConstantStr =
165+
auto *ConstantStr =
167166
dyn_cast<ConstantDataSequential>(GlobalStr->getOperand(0));
168167
if (!ConstantStr)
169168
report_fatal_error("__nvvm_reflect argument must be a string constant");
170169
if (!ConstantStr->isCString())
171170
report_fatal_error(
172171
"__nvvm_reflect argument must be a null-terminated string");
173172

174-
const StringRef ReflectArg = ConstantStr->getAsString().drop_back();
173+
StringRef ReflectArg = ConstantStr->getAsString().drop_back();
175174
if (ReflectArg.empty())
176175
report_fatal_error("__nvvm_reflect argument cannot be empty");
177176
// Now that we have extracted the string argument, we can look it up in the
@@ -183,7 +182,7 @@ bool NVVMReflect::handleReflectFunction(Module &M,
183182
LLVM_DEBUG(dbgs() << "Replacing call of reflect function " << F->getName()
184183
<< "(" << ReflectArg << ") with value " << ReflectVal
185184
<< "\n");
186-
auto *const NewValue = ConstantInt::get(Call->getType(), ReflectVal);
185+
auto *NewValue = ConstantInt::get(Call->getType(), ReflectVal);
187186
foldReflectCall(Call, NewValue);
188187
Call->eraseFromParent();
189188
}
@@ -193,25 +192,23 @@ bool NVVMReflect::handleReflectFunction(Module &M,
193192
return Changed;
194193
}
195194

196-
void NVVMReflect::foldReflectCall(CallInst *const Call,
197-
Constant *const NewValue) {
195+
void NVVMReflect::foldReflectCall(CallInst *Call, Constant *NewValue) {
198196
SmallVector<Instruction *, 8> Worklist;
199197
// Replace an instruction with a constant and add all users of the instruction
200198
// to the worklist
201-
auto ReplaceInstructionWithConst = [&](Instruction *const I,
202-
Constant *const C) {
203-
for (auto *const U : I->users())
204-
if (auto *const UI = dyn_cast<Instruction>(U))
199+
auto ReplaceInstructionWithConst = [&](Instruction *I, Constant *C) {
200+
for (auto *U : I->users())
201+
if (auto *UI = dyn_cast<Instruction>(U))
205202
Worklist.push_back(UI);
206203
I->replaceAllUsesWith(C);
207204
};
208205

209206
ReplaceInstructionWithConst(Call, NewValue);
210207

211-
const auto &DL = Call->getModule()->getDataLayout();
208+
auto &DL = Call->getModule()->getDataLayout();
212209
while (!Worklist.empty()) {
213-
auto *const I = Worklist.pop_back_val();
214-
if (auto *const C = ConstantFoldInstruction(I, DL)) {
210+
auto *I = Worklist.pop_back_val();
211+
if (auto *C = ConstantFoldInstruction(I, DL)) {
215212
ReplaceInstructionWithConst(I, C);
216213
if (isInstructionTriviallyDead(I))
217214
I->eraseFromParent();

0 commit comments

Comments
 (0)