Skip to content

Commit 2b44fbb

Browse files
SC llvm teamSC llvm team
SC llvm team
authored and
SC llvm team
committed
Merged main:ea72a4e6547f into amd-gfx:7199f6eb6d6f
Local branch amd-gfx 7199f6e Merged main:211c9752c820 into amd-gfx:1dd68fefc1fa Remote branch main ea72a4e [CUDA][HIP] Fix template argument deduction
2 parents 7199f6e + ea72a4e commit 2b44fbb

File tree

53 files changed

+716
-842
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

53 files changed

+716
-842
lines changed

‎clang-tools-extra/clang-tidy/readability/ImplicitBoolConversionCheck.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -228,7 +228,8 @@ bool isCastAllowedInCondition(const ImplicitCastExpr *Cast,
228228
if (!S)
229229
return false;
230230
if (isa<IfStmt>(S) || isa<ConditionalOperator>(S) || isa<ForStmt>(S) ||
231-
isa<WhileStmt>(S) || isa<BinaryConditionalOperator>(S))
231+
isa<WhileStmt>(S) || isa<DoStmt>(S) ||
232+
isa<BinaryConditionalOperator>(S))
232233
return true;
233234
if (isa<ParenExpr>(S) || isa<ImplicitCastExpr>(S) ||
234235
isUnaryLogicalNotOperator(S) ||

‎clang-tools-extra/docs/ReleaseNotes.rst

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -196,6 +196,11 @@ Changes in existing checks
196196
<clang-tidy/checks/readability/identifier-naming>` check to emit proper
197197
warnings when a type forward declaration precedes its definition.
198198

199+
- Improved :doc:`readability-implicit-bool-conversion
200+
<clang-tidy/checks/readability/implicit-bool-conversion>` check to take
201+
do-while loops into account for the `AllowIntegerConditions` and
202+
`AllowPointerConditions` options.
203+
199204
Removed checks
200205
^^^^^^^^^^^^^^
201206

‎clang-tools-extra/test/clang-tidy/checkers/readability/implicit-bool-conversion-allow-in-conditions.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,9 @@ void implicitConversionIntegerToBoolInConditionalsIsAllowed() {
3838
while (functionReturningInt()) {}
3939
while (functionReturningPointer()) {}
4040
while (functionReturningInt() && !functionReturningPointer() || (!functionReturningInt() && functionReturningPointer())) {}
41+
do {} while (functionReturningInt());
42+
do {} while (functionReturningPointer());
43+
do {} while (functionReturningInt() && !functionReturningPointer() || (!functionReturningInt() && functionReturningPointer()));
4144
int value1 = functionReturningInt() ? 1 : 2;
4245
int value2 = !functionReturningInt() ? 1 : 2;
4346
int value3 = (functionReturningInt() && functionReturningPointer() || !functionReturningInt()) ? 1 : 2;

‎clang/docs/ReleaseNotes.rst

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -123,6 +123,8 @@ Improvements to Clang's diagnostics
123123
template-specialization function calls.
124124
- Clang contexpr evaluator now displays notes as well as an error when a constructor
125125
of a base class is not called in the constructor of its derived class.
126+
- Clang no longer emits ``-Wmissing-variable-declarations`` for variables declared
127+
with the ``register`` storage class.
126128

127129
Bug Fixes in This Version
128130
-------------------------

‎clang/lib/Analysis/BodyFarm.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -806,7 +806,7 @@ static Stmt *createObjCPropertyGetter(ASTContext &Ctx,
806806

807807
if (!IVar) {
808808
Prop = MD->findPropertyDecl();
809-
IVar = findBackingIvar(Prop);
809+
IVar = Prop ? findBackingIvar(Prop) : nullptr;
810810
}
811811

812812
if (!IVar || !Prop)

‎clang/lib/Headers/hlsl/hlsl_intrinsics.h

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -532,5 +532,36 @@ __attribute__((clang_builtin_alias(__builtin_elementwise_bitreverse)))
532532
uint64_t3 reversebits(uint64_t3);
533533
__attribute__((clang_builtin_alias(__builtin_elementwise_bitreverse)))
534534
uint64_t4 reversebits(uint64_t4);
535+
536+
// pow builtins
537+
#ifdef __HLSL_ENABLE_16_BIT
538+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
539+
half pow(half, half);
540+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
541+
half2 pow(half2, half2);
542+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
543+
half3 pow(half3, half3);
544+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
545+
half4 pow(half4, half4);
546+
#endif
547+
548+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow))) float
549+
pow(float, float);
550+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
551+
float2 pow(float2, float2);
552+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
553+
float3 pow(float3, float3);
554+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
555+
float4 pow(float4, float4);
556+
557+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow))) double
558+
pow(double, double);
559+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
560+
double2 pow(double2, double2);
561+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
562+
double3 pow(double3, double3);
563+
__attribute__((clang_builtin_alias(__builtin_elementwise_pow)))
564+
double4 pow(double4, double4);
565+
535566
} // namespace hlsl
536567
#endif //_HLSL_HLSL_INTRINSICS_H_

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14144,6 +14144,7 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {
1414414144
var->getDeclContext()->getRedeclContext()->isFileContext() &&
1414514145
var->isExternallyVisible() && var->hasLinkage() &&
1414614146
!var->isInline() && !var->getDescribedVarTemplate() &&
14147+
var->getStorageClass() != SC_Register &&
1414714148
!isa<VarTemplatePartialSpecializationDecl>(var) &&
1414814149
!isTemplateInstantiation(var->getTemplateSpecializationKind()) &&
1414914150
!getDiagnostics().isIgnored(diag::warn_missing_variable_declarations,

‎clang/lib/Sema/SemaOverload.cpp

Lines changed: 37 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -12770,6 +12770,13 @@ Sema::resolveAddressOfSingleOverloadCandidate(Expr *E, DeclAccessPair &Pair) {
1277012770
DeclAccessPair DAP;
1277112771
SmallVector<FunctionDecl *, 2> AmbiguousDecls;
1277212772

12773+
// Return positive for better, negative for worse, 0 for equal preference.
12774+
auto CheckCUDAPreference = [&](FunctionDecl *FD1, FunctionDecl *FD2) {
12775+
FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
12776+
return static_cast<int>(IdentifyCUDAPreference(Caller, FD1)) -
12777+
static_cast<int>(IdentifyCUDAPreference(Caller, FD2));
12778+
};
12779+
1277312780
auto CheckMoreConstrained = [&](FunctionDecl *FD1,
1277412781
FunctionDecl *FD2) -> std::optional<bool> {
1277512782
if (FunctionDecl *MF = FD1->getInstantiatedFromMemberFunction())
@@ -12800,9 +12807,31 @@ Sema::resolveAddressOfSingleOverloadCandidate(Expr *E, DeclAccessPair &Pair) {
1280012807
if (!checkAddressOfFunctionIsAvailable(FD))
1280112808
continue;
1280212809

12810+
// If we found a better result, update Result.
12811+
auto FoundBetter = [&]() {
12812+
IsResultAmbiguous = false;
12813+
DAP = I.getPair();
12814+
Result = FD;
12815+
};
12816+
1280312817
// We have more than one result - see if it is more constrained than the
1280412818
// previous one.
1280512819
if (Result) {
12820+
// Check CUDA preference first. If the candidates have differennt CUDA
12821+
// preference, choose the one with higher CUDA preference. Otherwise,
12822+
// choose the one with more constraints.
12823+
if (getLangOpts().CUDA) {
12824+
int PreferenceByCUDA = CheckCUDAPreference(FD, Result);
12825+
// FD has different preference than Result.
12826+
if (PreferenceByCUDA != 0) {
12827+
// FD is more preferable than Result.
12828+
if (PreferenceByCUDA > 0)
12829+
FoundBetter();
12830+
continue;
12831+
}
12832+
}
12833+
// FD has the same CUDA prefernece than Result. Continue check
12834+
// constraints.
1280612835
std::optional<bool> MoreConstrainedThanPrevious =
1280712836
CheckMoreConstrained(FD, Result);
1280812837
if (!MoreConstrainedThanPrevious) {
@@ -12814,9 +12843,7 @@ Sema::resolveAddressOfSingleOverloadCandidate(Expr *E, DeclAccessPair &Pair) {
1281412843
continue;
1281512844
// FD is more constrained - replace Result with it.
1281612845
}
12817-
IsResultAmbiguous = false;
12818-
DAP = I.getPair();
12819-
Result = FD;
12846+
FoundBetter();
1282012847
}
1282112848

1282212849
if (IsResultAmbiguous)
@@ -12826,9 +12853,15 @@ Sema::resolveAddressOfSingleOverloadCandidate(Expr *E, DeclAccessPair &Pair) {
1282612853
SmallVector<const Expr *, 1> ResultAC;
1282712854
// We skipped over some ambiguous declarations which might be ambiguous with
1282812855
// the selected result.
12829-
for (FunctionDecl *Skipped : AmbiguousDecls)
12856+
for (FunctionDecl *Skipped : AmbiguousDecls) {
12857+
// If skipped candidate has different CUDA preference than the result,
12858+
// there is no ambiguity. Otherwise check whether they have different
12859+
// constraints.
12860+
if (getLangOpts().CUDA && CheckCUDAPreference(Skipped, Result) != 0)
12861+
continue;
1283012862
if (!CheckMoreConstrained(Skipped, Result))
1283112863
return nullptr;
12864+
}
1283212865
Pair = DAP;
1283312866
}
1283412867
return Result;

‎clang/test/CodeGenHLSL/builtins/pow.hlsl

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
// RUN: %clang_cc1 -std=hlsl2021 -finclude-default-header -x hlsl -triple \
2+
// RUN: dxil-pc-shadermodel6.3-library %s -fnative-half-type \
3+
// RUN: -emit-llvm -disable-llvm-passes -O3 -o - | FileCheck %s
4+
// RUN: %clang_cc1 -std=hlsl2021 -finclude-default-header -x hlsl -triple \
5+
// RUN: dxil-pc-shadermodel6.3-library %s -emit-llvm -disable-llvm-passes \
6+
// RUN: -D__HLSL_ENABLE_16_BIT -o - | FileCheck %s --check-prefix=NO_HALF
7+
8+
// CHECK: define noundef half @
9+
// CHECK: call half @llvm.pow.f16(
10+
// NO_HALF: define noundef float @"?test_pow_half@@YA$halff@$halff@0@Z"(
11+
// NO_HALF: call float @llvm.pow.f32(
12+
half test_pow_half(half p0, half p1)
13+
{
14+
return pow(p0, p1);
15+
}
16+
// CHECK: define noundef <2 x half> @"?test_pow_half2@@YAT?$__vector@$f16@$01@__clang@@T12@0@Z"(
17+
// CHECK: call <2 x half> @llvm.pow.v2f16
18+
// NO_HALF: define noundef <2 x float> @"?test_pow_float2@@YAT?$__vector@M$01@__clang@@T12@0@Z"(
19+
// NO_HALF: call <2 x float> @llvm.pow.v2f32(
20+
half2 test_pow_half2(half2 p0, half2 p1)
21+
{
22+
return pow(p0, p1);
23+
}
24+
// CHECK: define noundef <3 x half> @"?test_pow_half3@@YAT?$__vector@$f16@$02@__clang@@T12@0@Z"(
25+
// CHECK: call <3 x half> @llvm.pow.v3f16
26+
// NO_HALF: define noundef <3 x float> @"?test_pow_float3@@YAT?$__vector@M$02@__clang@@T12@0@Z"(
27+
// NO_HALF: call <3 x float> @llvm.pow.v3f32(
28+
half3 test_pow_half3(half3 p0, half3 p1)
29+
{
30+
return pow(p0, p1);
31+
}
32+
// CHECK: define noundef <4 x half> @"?test_pow_half4@@YAT?$__vector@$f16@$03@__clang@@T12@0@Z"(
33+
// CHECK: call <4 x half> @llvm.pow.v4f16
34+
// NO_HALF: define noundef <4 x float> @"?test_pow_float4@@YAT?$__vector@M$03@__clang@@T12@0@Z"(
35+
// NO_HALF: call <4 x float> @llvm.pow.v4f32(
36+
half4 test_pow_half4(half4 p0, half4 p1)
37+
{
38+
return pow(p0, p1);
39+
}
40+
41+
// CHECK: define noundef float @"?test_pow_float@@YAMMM@Z"(
42+
// CHECK: call float @llvm.pow.f32(
43+
float test_pow_float(float p0, float p1)
44+
{
45+
return pow(p0, p1);
46+
}
47+
// CHECK: define noundef <2 x float> @"?test_pow_float2@@YAT?$__vector@M$01@__clang@@T12@0@Z"(
48+
// CHECK: call <2 x float> @llvm.pow.v2f32
49+
float2 test_pow_float2(float2 p0, float2 p1)
50+
{
51+
return pow(p0, p1);
52+
}
53+
// CHECK: define noundef <3 x float> @"?test_pow_float3@@YAT?$__vector@M$02@__clang@@T12@0@Z"(
54+
// CHECK: call <3 x float> @llvm.pow.v3f32
55+
float3 test_pow_float3(float3 p0, float3 p1)
56+
{
57+
return pow(p0, p1);
58+
}
59+
// CHECK: define noundef <4 x float> @"?test_pow_float4@@YAT?$__vector@M$03@__clang@@T12@0@Z"(
60+
// CHECK: call <4 x float> @llvm.pow.v4f32
61+
float4 test_pow_float4(float4 p0, float4 p1)
62+
{
63+
return pow(p0, p1);
64+
}
65+
66+
// CHECK: define noundef double @"?test_pow_double@@YANNN@Z"(
67+
// CHECK: call double @llvm.pow.f64(
68+
double test_pow_double(double p0, double p1)
69+
{
70+
return pow(p0, p1);
71+
}
72+
// CHECK: define noundef <2 x double> @"?test_pow_double2@@YAT?$__vector@N$01@__clang@@T12@0@Z"(
73+
// CHECK: call <2 x double> @llvm.pow.v2f64
74+
double2 test_pow_double2(double2 p0, double2 p1)
75+
{
76+
return pow(p0, p1);
77+
}
78+
// CHECK: define noundef <3 x double> @"?test_pow_double3@@YAT?$__vector@N$02@__clang@@T12@0@Z"(
79+
// CHECK: call <3 x double> @llvm.pow.v3f64
80+
double3 test_pow_double3(double3 p0, double3 p1)
81+
{
82+
return pow(p0, p1);
83+
}
84+
// CHECK: define noundef <4 x double> @"?test_pow_double4@@YAT?$__vector@N$03@__clang@@T12@0@Z"(
85+
// CHECK: call <4 x double> @llvm.pow.v4f64
86+
double4 test_pow_double4(double4 p0, double4 p1)
87+
{
88+
return pow(p0, p1);
89+
}

‎clang/test/Sema/warn-missing-variable-declarations-register.c

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,4 @@
1+
// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -Wmissing-variable-declarations -fsyntax-only -verify %s
2+
// expected-no-diagnostics
3+
4+
register unsigned long current_stack_pointer asm("rsp");

‎clang/test/SemaCUDA/template-arg-deduction.cu

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
2+
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
3+
4+
// expected-no-diagnostics
5+
6+
#include "Inputs/cuda.h"
7+
8+
void foo();
9+
__device__ void foo();
10+
11+
template<class F>
12+
void host_temp(F f);
13+
14+
template<class F>
15+
__device__ void device_temp(F f);
16+
17+
void host_caller() {
18+
host_temp(foo);
19+
}
20+
21+
__global__ void kernel_caller() {
22+
device_temp(foo);
23+
}
24+
25+
__device__ void device_caller() {
26+
device_temp(foo);
27+
}

‎llvm/bindings/ocaml/llvm/META.llvm.in

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ package "bitwriter" (
3030
)
3131

3232
package "executionengine" (
33-
requires = "llvm,llvm.target,ctypes.foreign"
33+
requires = "llvm,llvm.target,ctypes"
3434
version = "@PACKAGE_VERSION@"
3535
description = "JIT and Interpreter for LLVM"
3636
archive(byte) = "llvm_executionengine.cma"

‎llvm/include/llvm/Analysis/ValueTracking.h

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -124,10 +124,6 @@ bool isKnownToBeAPowerOfTwo(const Value *V, const DataLayout &DL,
124124
const DominatorTree *DT = nullptr,
125125
bool UseInstrInfo = true);
126126

127-
/// Return true if the given instruction is only used in zero comparison
128-
bool isOnlyUsedInZeroComparison(const Instruction *CxtI);
129-
130-
/// Return true if the given instruction is only used in zero equality comparison
131127
bool isOnlyUsedInZeroEqualityComparison(const Instruction *CxtI);
132128

133129
/// Return true if the given value is known to be non-zero when defined. For

‎llvm/include/llvm/Config/llvm-config.h.cmake

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616

1717
/* Indicate that this is LLVM compiled from the amd-gfx branch. */
1818
#define LLVM_HAVE_BRANCH_AMD_GFX
19-
#define LLVM_MAIN_REVISION 470569
19+
#define LLVM_MAIN_REVISION 470586
2020

2121
/* Define if LLVM_ENABLE_DUMP is enabled */
2222
#cmakedefine LLVM_ENABLE_DUMP

‎llvm/include/llvm/ExecutionEngine/Orc/Core.h

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1434,6 +1434,9 @@ class ExecutionSession {
14341434
/// Return the triple for the executor.
14351435
const Triple &getTargetTriple() const { return EPC->getTargetTriple(); }
14361436

1437+
// Return the page size for the executor.
1438+
size_t getPageSize() const { return EPC->getPageSize(); }
1439+
14371440
/// Get the SymbolStringPool for this instance.
14381441
std::shared_ptr<SymbolStringPool> getSymbolStringPool() {
14391442
return EPC->getSymbolStringPool();

‎llvm/include/llvm/Transforms/AggressiveInstCombine/AggressiveInstCombine.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
/// \file
99
///
1010
/// AggressiveInstCombiner - Combine expression patterns to form expressions
11-
/// with fewer, simple instructions.
11+
/// with fewer, simple instructions. This pass does not modify the CFG.
1212
///
1313
//===----------------------------------------------------------------------===//
1414

‎llvm/lib/Analysis/ValueTracking.cpp

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -261,13 +261,6 @@ bool llvm::haveNoCommonBitsSet(const Value *LHS, const Value *RHS,
261261
return KnownBits::haveNoCommonBitsSet(LHSKnown, RHSKnown);
262262
}
263263

264-
bool llvm::isOnlyUsedInZeroComparison(const Instruction *I) {
265-
return !I->user_empty() && all_of(I->users(), [](const User *U) {
266-
ICmpInst::Predicate P;
267-
return match(U, m_ICmp(P, m_Value(), m_Zero()));
268-
});
269-
}
270-
271264
bool llvm::isOnlyUsedInZeroEqualityComparison(const Instruction *I) {
272265
return !I->user_empty() && all_of(I->users(), [](const User *U) {
273266
ICmpInst::Predicate P;

‎llvm/lib/Target/AArch64/AArch64ISelLowering.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2107,6 +2107,13 @@ void AArch64TargetLowering::computeKnownBitsForTargetNode(
21072107
Known = KnownBits::ashr(Known, Known2);
21082108
break;
21092109
}
2110+
case AArch64ISD::VSHL: {
2111+
KnownBits Known2;
2112+
Known = DAG.computeKnownBits(Op->getOperand(0), Depth + 1);
2113+
Known2 = DAG.computeKnownBits(Op->getOperand(1), Depth + 1);
2114+
Known = KnownBits::shl(Known, Known2);
2115+
break;
2116+
}
21102117
case AArch64ISD::MOVI: {
21112118
ConstantSDNode *CN = cast<ConstantSDNode>(Op->getOperand(0));
21122119
Known =

0 commit comments

Comments
 (0)