Skip to content

Commit 98ccf24

Browse files
tonykuttaitonykuttai
authored
Merge branch 'main' into himadhith/xxleqv_vec
2 parents 39431f2 + cc5185b commit 98ccf24

22 files changed

+2575
-49
lines changed

‎clang/lib/Sema/SemaModule.cpp‎

Lines changed: 20 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1282,7 +1282,16 @@ bool ExposureChecker::isExposureCandidate(const NamedDecl *D) {
12821282
// (outside the private-module-fragment, if any) or
12831283
// module partition is an exposure, the program is ill-formed.
12841284
Module *M = D->getOwningModule();
1285-
if (!M || !M->isInterfaceOrPartition())
1285+
if (!M)
1286+
return false;
1287+
// If M is implicit global module, the declaration must be in the purview of
1288+
// a module unit.
1289+
if (M->isImplicitGlobalModule()) {
1290+
M = M->Parent;
1291+
assert(M && "Implicit global module must have a parent");
1292+
}
1293+
1294+
if (!M->isInterfaceOrPartition())
12861295
return false;
12871296

12881297
if (D->isImplicit())
@@ -1495,6 +1504,16 @@ bool ExposureChecker::checkExposure(const Stmt *S, bool Diag) {
14951504

14961505
void ExposureChecker::checkExposureInContext(const DeclContext *DC) {
14971506
for (auto *TopD : DC->noload_decls()) {
1507+
if (auto *Export = dyn_cast<ExportDecl>(TopD)) {
1508+
checkExposureInContext(Export);
1509+
continue;
1510+
}
1511+
1512+
if (auto *LinkageSpec = dyn_cast<LinkageSpecDecl>(TopD)) {
1513+
checkExposureInContext(LinkageSpec);
1514+
continue;
1515+
}
1516+
14981517
auto *TopND = dyn_cast<NamedDecl>(TopD);
14991518
if (!TopND)
15001519
continue;

‎clang/lib/Sema/SemaOpenMP.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1364,15 +1364,15 @@ DSAStackTy::DSAVarData DSAStackTy::getDSA(const_iterator &Iter,
13641364
DefaultDataSharingAttributes IterDA = Iter->DefaultAttr;
13651365
switch (Iter->DefaultVCAttr) {
13661366
case DSA_VC_aggregate:
1367-
if (!VD->getType()->isAggregateType())
1367+
if (!D->getType()->isAggregateType())
13681368
IterDA = DSA_none;
13691369
break;
13701370
case DSA_VC_pointer:
1371-
if (!VD->getType()->isPointerType())
1371+
if (!D->getType()->isPointerType())
13721372
IterDA = DSA_none;
13731373
break;
13741374
case DSA_VC_scalar:
1375-
if (!VD->getType()->isScalarType())
1375+
if (!D->getType()->isScalarType())
13761376
IterDA = DSA_none;
13771377
break;
13781378
case DSA_VC_all:

‎clang/test/Modules/reference-tu-local-var.cppm‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,9 @@
1+
// RUN: %clang_cc1 -std=c++23 %s -verify -fsyntax-only
2+
export module M;
3+
static int local;
4+
export inline int exposure1() { return local; } // expected-warning {{TU local entity 'local' is exposed}}
5+
6+
static int local2 = 43;
7+
export extern "C++" {
8+
inline int exposure2() { return local2; } // expected-warning {{TU local entity 'local2' is exposed}}
9+
}

‎clang/test/OpenMP/parallel_default_variableCategory_codegen.cpp‎

Lines changed: 86 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,86 @@
1+
// RUN: %clangxx -Xclang -verify -Wno-vla -fopenmp -fopenmp-version=60 -x c++ -S -emit-llvm %s -o - | FileCheck %s
2+
// expected-no-diagnostics
3+
#ifndef HEADER
4+
#define HEADER
5+
6+
int global;
7+
#define VECTOR_SIZE 4
8+
9+
int main (int argc, char **argv) {
10+
int i,n;
11+
int x;
12+
13+
n = VECTOR_SIZE;
14+
15+
#pragma omp parallel masked firstprivate(x) num_threads(2)
16+
{
17+
int *xPtr = nullptr;
18+
// scalar
19+
#pragma omp task default(shared:scalar)
20+
{
21+
xPtr = &x;
22+
}
23+
#pragma omp taskwait
24+
25+
// pointer
26+
#pragma omp task default(shared:pointer) shared(x)
27+
{
28+
xPtr = &x;
29+
}
30+
#pragma omp taskwait
31+
}
32+
33+
int *aggregate[VECTOR_SIZE] = {0,0,0,0};
34+
35+
#pragma omp parallel masked num_threads(2)
36+
{
37+
// aggregate
38+
#pragma omp task default(shared:aggregate)
39+
for(i=0;i<n;i++) {
40+
aggregate[i] = &x;
41+
}
42+
#pragma omp taskwait
43+
44+
#pragma omp task default(shared:aggregate) shared(x)
45+
for(i=0;i<n;i++) {
46+
aggregate[i] = &x;
47+
}
48+
#pragma omp taskwait
49+
50+
// all
51+
#pragma omp task default(shared:all)
52+
for(i=0;i<n;i++) {
53+
aggregate[i] = &x;
54+
}
55+
#pragma omp taskwait
56+
}
57+
}
58+
59+
#endif
60+
61+
// CHECK-LABEL: define {{.*}}main.omp_outlined{{.*}}
62+
// CHECK: store ptr null, ptr{{.*}}
63+
// CHECK-NEXT: {{.*}}getelementptr {{.*}}
64+
// CHECK-NEXT: store ptr {{.*}}
65+
// CHECK-NEXT: {{.*}}getelementptr {{.*}}
66+
// CHECK-NEXT: store ptr {{.*}}
67+
// CHECK-NEXT: {{.*}}call{{.*}}__kmpc_omp_task_alloc{{.*}}
68+
// CHECK: ret void
69+
//
70+
// CHECK: define {{.*}}main.omp_outlined{{.*}}
71+
// CHECK: {{.*}}getelementptr {{.*}}
72+
// CHECK-NEXT: store ptr {{.*}}
73+
// CHECK-NEXT: {{.*}}call{{.*}}__kmpc_omp_task_alloc{{.*}}
74+
// CHECK: store ptr {{.*}}
75+
// CHECK-NEXT: {{.*}}getelementptr {{.*}}
76+
// CHECK-NEXT: store ptr {{.*}}
77+
// CHECK-NEXT: {{.*}}call{{.*}}__kmpc_omp_task_alloc{{.*}}
78+
// CHECK: store ptr {{.*}}
79+
// CHECK-NEXT: {{.*}}getelementptr {{.*}}
80+
// CHECK-NEXT: store ptr {{.*}}
81+
// CHECK-NEXT: {{.*}}getelementptr {{.*}}
82+
// CHECK-NEXT: store ptr {{.*}}
83+
// CHECK-NEXT: {{.*}}getelementptr {{.*}}
84+
// CHECK-NEXT: store ptr {{.*}}
85+
// CHECK-NEXT: {{.*}}call{{.*}}__kmpc_omp_task_alloc{{.*}}
86+
// CHECK: ret void

‎lldb/include/lldb/Utility/NonNullSharedPtr.h‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -29,12 +29,14 @@ template <typename T> class NonNullSharedPtr : private std::shared_ptr<T> {
2929
public:
3030
NonNullSharedPtr(const std::shared_ptr<T> &t)
3131
: Base(t ? t : std::make_shared<T>()) {
32-
assert(t && "NonNullSharedPtr initialized from NULL shared_ptr");
32+
assert(t && "NonNullSharedPtr constructed from nullptr");
3333
}
3434

35-
NonNullSharedPtr(std::shared_ptr<T> &&t)
36-
: Base(t ? std::move(t) : std::make_shared<T>()) {
37-
// Can't assert on t as it's been moved-from.
35+
NonNullSharedPtr(std::shared_ptr<T> &&t) : Base(std::move(t)) {
36+
const auto b = static_cast<bool>(*this);
37+
assert(b && "NonNullSharedPtr constructed from nullptr");
38+
if (!b)
39+
Base::operator=(std::make_shared<T>());
3840
}
3941

4042
NonNullSharedPtr(const NonNullSharedPtr &other) : Base(other) {}

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td‎

Lines changed: 30 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -2071,34 +2071,36 @@ def : Pat<(int_nvvm_ull2d_rp i64:$a), (CVT_f64_u64 $a, CvtRP)>;
20712071
def : Pat<(int_nvvm_f2h_rn_ftz f32:$a), (CVT_f16_f32 $a, CvtRN_FTZ)>;
20722072
def : Pat<(int_nvvm_f2h_rn f32:$a), (CVT_f16_f32 $a, CvtRN)>;
20732073

2074-
def : Pat<(int_nvvm_ff_to_e4m3x2_rn f32:$a, f32:$b),
2075-
(CVT_e4m3x2_f32 $a, $b, CvtRN)>;
2076-
def : Pat<(int_nvvm_ff_to_e4m3x2_rn_relu f32:$a, f32:$b),
2077-
(CVT_e4m3x2_f32 $a, $b, CvtRN_RELU)>;
2078-
def : Pat<(int_nvvm_ff_to_e5m2x2_rn f32:$a, f32:$b),
2079-
(CVT_e5m2x2_f32 $a, $b, CvtRN)>;
2080-
def : Pat<(int_nvvm_ff_to_e5m2x2_rn_relu f32:$a, f32:$b),
2081-
(CVT_e5m2x2_f32 $a, $b, CvtRN_RELU)>;
2082-
2083-
def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn v2f16:$a),
2084-
(CVT_e4m3x2_f16x2 $a, CvtRN)>;
2085-
def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn_relu v2f16:$a),
2086-
(CVT_e4m3x2_f16x2 $a, CvtRN_RELU)>;
2087-
def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn v2f16:$a),
2088-
(CVT_e5m2x2_f16x2 $a, CvtRN)>;
2089-
def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn_relu v2f16:$a),
2090-
(CVT_e5m2x2_f16x2 $a, CvtRN_RELU)>;
2091-
2092-
def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn i16:$a),
2093-
(CVT_f16x2_e4m3x2 $a, CvtRN)>;
2094-
def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn_relu i16:$a),
2095-
(CVT_f16x2_e4m3x2 $a, CvtRN_RELU)>;
2096-
def : Pat<(int_nvvm_e5m2x2_to_f16x2_rn i16:$a),
2097-
(CVT_f16x2_e5m2x2 $a, CvtRN)>;
2098-
def : Pat<(int_nvvm_e5m2x2_to_f16x2_rn_relu i16:$a),
2099-
(CVT_f16x2_e5m2x2 $a, CvtRN_RELU)>;
2100-
2101-
let Predicates = [hasPTX<86>, hasSM<100>, hasArchAccelFeatures] in {
2074+
let Predicates = [callSubtarget<"hasFP8ConversionSupport">] in {
2075+
def : Pat<(int_nvvm_ff_to_e4m3x2_rn f32:$a, f32:$b),
2076+
(CVT_e4m3x2_f32 $a, $b, CvtRN)>;
2077+
def : Pat<(int_nvvm_ff_to_e4m3x2_rn_relu f32:$a, f32:$b),
2078+
(CVT_e4m3x2_f32 $a, $b, CvtRN_RELU)>;
2079+
def : Pat<(int_nvvm_ff_to_e5m2x2_rn f32:$a, f32:$b),
2080+
(CVT_e5m2x2_f32 $a, $b, CvtRN)>;
2081+
def : Pat<(int_nvvm_ff_to_e5m2x2_rn_relu f32:$a, f32:$b),
2082+
(CVT_e5m2x2_f32 $a, $b, CvtRN_RELU)>;
2083+
2084+
def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn v2f16:$a),
2085+
(CVT_e4m3x2_f16x2 $a, CvtRN)>;
2086+
def : Pat<(int_nvvm_f16x2_to_e4m3x2_rn_relu v2f16:$a),
2087+
(CVT_e4m3x2_f16x2 $a, CvtRN_RELU)>;
2088+
def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn v2f16:$a),
2089+
(CVT_e5m2x2_f16x2 $a, CvtRN)>;
2090+
def : Pat<(int_nvvm_f16x2_to_e5m2x2_rn_relu v2f16:$a),
2091+
(CVT_e5m2x2_f16x2 $a, CvtRN_RELU)>;
2092+
2093+
def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn i16:$a),
2094+
(CVT_f16x2_e4m3x2 $a, CvtRN)>;
2095+
def : Pat<(int_nvvm_e4m3x2_to_f16x2_rn_relu i16:$a),
2096+
(CVT_f16x2_e4m3x2 $a, CvtRN_RELU)>;
2097+
def : Pat<(int_nvvm_e5m2x2_to_f16x2_rn i16:$a),
2098+
(CVT_f16x2_e5m2x2 $a, CvtRN)>;
2099+
def : Pat<(int_nvvm_e5m2x2_to_f16x2_rn_relu i16:$a),
2100+
(CVT_f16x2_e5m2x2 $a, CvtRN_RELU)>;
2101+
}
2102+
2103+
let Predicates = [callSubtarget<"hasNarrowFPConversionSupport">] in {
21022104
def : Pat<(int_nvvm_ff_to_e2m3x2_rn_satfinite f32:$a, f32:$b),
21032105
(CVT_e2m3x2_f32_sf $a, $b, CvtRN)>;
21042106
def : Pat<(int_nvvm_ff_to_e2m3x2_rn_relu_satfinite f32:$a, f32:$b),

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

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -177,6 +177,27 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
177177
hasPTXWithAccelSMs(86, {100, 101});
178178
}
179179

180+
// Checks support for conversions involving e4m3x2 and e5m2x2.
181+
bool hasFP8ConversionSupport() const {
182+
if (PTXVersion >= 81)
183+
return SmVersion >= 89;
184+
185+
if (PTXVersion >= 78)
186+
return SmVersion >= 90;
187+
188+
return false;
189+
}
190+
191+
// Checks support for conversions involving the following types:
192+
// - e2m3x2/e3m2x2
193+
// - e2m1x2
194+
// - ue8m0x2
195+
bool hasNarrowFPConversionSupport() const {
196+
return hasPTXWithFamilySMs(90, {100, 110, 120}) ||
197+
hasPTXWithFamilySMs(88, {100, 101, 120}) ||
198+
hasPTXWithAccelSMs(86, {100, 101, 120});
199+
}
200+
180201
// Prior to CUDA 12.3 ptxas did not recognize that the trap instruction
181202
// terminates a basic block. Instead, it would assume that control flow
182203
// continued to the next instruction. The next instruction could be in the

‎llvm/lib/Target/PowerPC/PPCISelLowering.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15563,8 +15563,8 @@ SDValue convertTwoLoadsAndCmpToVCMPEQUB(SelectionDAG &DAG, SDNode *N,
1556315563
assert(N->getOpcode() == ISD::SETCC && "Should be called with a SETCC node");
1556415564

1556515565
ISD::CondCode CC = cast<CondCodeSDNode>(N->getOperand(2))->get();
15566-
assert(CC == ISD::SETNE ||
15567-
CC == ISD::SETEQ && "CC mus be ISD::SETNE or ISD::SETEQ");
15566+
assert((CC == ISD::SETNE || CC == ISD::SETEQ) &&
15567+
"CC mus be ISD::SETNE or ISD::SETEQ");
1556815568

1556915569
auto getV16i8Load = [&](const SDValue &Operand) {
1557015570
if (Operand.getOpcode() == ISD::Constant)

‎llvm/lib/Target/RISCV/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,7 @@ add_llvm_target(RISCVCodeGen
7272
RISCVVLOptimizer.cpp
7373
RISCVVMV0Elimination.cpp
7474
RISCVZacasABIFix.cpp
75+
RISCVZilsdOptimizer.cpp
7576
GISel/RISCVCallLowering.cpp
7677
GISel/RISCVInstructionSelector.cpp
7778
GISel/RISCVLegalizerInfo.cpp

‎llvm/lib/Target/RISCV/RISCV.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,6 +102,9 @@ void initializeRISCVPushPopOptPass(PassRegistry &);
102102
FunctionPass *createRISCVLoadStoreOptPass();
103103
void initializeRISCVLoadStoreOptPass(PassRegistry &);
104104

105+
FunctionPass *createRISCVPreAllocZilsdOptPass();
106+
void initializeRISCVPreAllocZilsdOptPass(PassRegistry &);
107+
105108
FunctionPass *createRISCVZacasABIFixPass();
106109
void initializeRISCVZacasABIFixPass(PassRegistry &);
107110

0 commit comments

Comments
 (0)