Skip to content

Commit f11aa52

Browse files
SC llvm teamSC llvm team
SC llvm team
authored and
SC llvm team
committed
Merged main:e04dd68a3a26d3ebdc2db07cf2f8807a02d30ce2 into amd-gfx:26b5e3593208
Local branch amd-gfx 26b5e35 Merged main:5373daad9492e157c0c1ad496334f5dfd78d7da0 into amd-gfx:0ab8eb800a42 Remote branch main e04dd68 [GlobalMerge] Use vector::assign in place of fill+resize. NFC (llvm#85723)
2 parents 26b5e35 + e04dd68 commit f11aa52

File tree

21 files changed

+381
-208
lines changed

21 files changed

+381
-208
lines changed

‎clang/lib/Driver/ToolChains/WebAssembly.cpp

Lines changed: 25 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -44,8 +44,15 @@ std::string wasm::Linker::getLinkerPath(const ArgList &Args) const {
4444
llvm::sys::fs::can_execute(UseLinker))
4545
return std::string(UseLinker);
4646

47-
// Accept 'lld', and 'ld' as aliases for the default linker
48-
if (UseLinker != "lld" && UseLinker != "ld")
47+
// Interpret 'lld' as explicitly requesting `wasm-ld`, so look for that
48+
// linker. Note that for `wasm32-wasip2` this overrides the default linker
49+
// of `wasm-component-ld`.
50+
if (UseLinker == "lld") {
51+
return ToolChain.GetProgramPath("wasm-ld");
52+
}
53+
54+
// Allow 'ld' as an alias for the default linker
55+
if (UseLinker != "ld")
4956
ToolChain.getDriver().Diag(diag::err_drv_invalid_linker_name)
5057
<< A->getAsString(Args);
5158
}
@@ -73,6 +80,16 @@ void wasm::Linker::ConstructJob(Compilation &C, const JobAction &JA,
7380
if (Args.hasArg(options::OPT_s))
7481
CmdArgs.push_back("--strip-all");
7582

83+
// On `wasip2` the default linker is `wasm-component-ld` which wraps the
84+
// execution of `wasm-ld`. Find `wasm-ld` and pass it as an argument of where
85+
// to find it to avoid it needing to hunt and rediscover or search `PATH` for
86+
// where it is.
87+
if (llvm::sys::path::stem(Linker).ends_with_insensitive(
88+
"wasm-component-ld")) {
89+
CmdArgs.push_back("--wasm-ld-path");
90+
CmdArgs.push_back(Args.MakeArgString(ToolChain.GetProgramPath("wasm-ld")));
91+
}
92+
7693
Args.addAllArgs(CmdArgs, {options::OPT_L, options::OPT_u});
7794

7895
ToolChain.AddFilePathLibArgs(Args, CmdArgs);
@@ -221,6 +238,12 @@ WebAssembly::WebAssembly(const Driver &D, const llvm::Triple &Triple,
221238
}
222239
}
223240

241+
const char *WebAssembly::getDefaultLinker() const {
242+
if (getOS() == "wasip2")
243+
return "wasm-component-ld";
244+
return "wasm-ld";
245+
}
246+
224247
bool WebAssembly::IsMathErrnoDefault() const { return false; }
225248

226249
bool WebAssembly::IsObjCNonFragileABIDefault() const { return true; }

‎clang/lib/Driver/ToolChains/WebAssembly.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,7 @@ class LLVM_LIBRARY_VISIBILITY WebAssembly final : public ToolChain {
6767
llvm::opt::ArgStringList &CmdArgs) const override;
6868
SanitizerMask getSupportedSanitizers() const override;
6969

70-
const char *getDefaultLinker() const override { return "wasm-ld"; }
70+
const char *getDefaultLinker() const override;
7171

7272
CXXStdlibType GetDefaultCXXStdlibType() const override {
7373
return ToolChain::CST_Libcxx;

‎clang/test/Driver/wasm-toolchain.c

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -197,3 +197,27 @@
197197
// RUN: not %clang -### %s --target=wasm32-unknown-unknown --sysroot=%s/no-sysroot-there -fPIC -mno-mutable-globals %s 2>&1 \
198198
// RUN: | FileCheck -check-prefix=PIC_NO_MUTABLE_GLOBALS %s
199199
// PIC_NO_MUTABLE_GLOBALS: error: invalid argument '-fPIC' not allowed with '-mno-mutable-globals'
200+
201+
// Test that `wasm32-wasip2` invokes the `wasm-component-ld` linker by default
202+
// instead of `wasm-ld`.
203+
204+
// RUN: %clang -### -O2 --target=wasm32-wasip2 %s --sysroot /foo 2>&1 \
205+
// RUN: | FileCheck -check-prefix=LINK_WASIP2 %s
206+
// LINK_WASIP2: "-cc1" {{.*}} "-o" "[[temp:[^"]*]]"
207+
// LINK_WASIP2: wasm-component-ld{{.*}}" "-L/foo/lib/wasm32-wasip2" "crt1.o" "[[temp]]" "-lc" "{{.*[/\\]}}libclang_rt.builtins-wasm32.a" "-o" "a.out"
208+
209+
// Test that on `wasm32-wasip2` the `wasm-component-ld` programs is told where
210+
// to find `wasm-ld` by default.
211+
212+
// RUN: %clang -### -O2 --target=wasm32-wasip2 %s --sysroot /foo 2>&1 \
213+
// RUN: | FileCheck -check-prefix=LINK_WASIP2_FIND_WASMLD %s
214+
// LINK_WASIP2_FIND_WASMLD: "-cc1" {{.*}} "-o" "[[temp:[^"]*]]"
215+
// LINK_WASIP2_FIND_WASMLD: wasm-component-ld{{.*}}" {{.*}} "--wasm-ld-path" "{{.*}}wasm-ld{{.*}}" {{.*}} "[[temp]]" {{.*}}
216+
217+
// If `wasm32-wasip2` is configured with `wasm-ld` as a linker then don't pass
218+
// the `--wasm-ld-path` flag.
219+
220+
// RUN: %clang -### -O2 --target=wasm32-wasip2 -fuse-ld=lld %s --sysroot /foo 2>&1 \
221+
// RUN: | FileCheck -check-prefix=LINK_WASIP2_USE_WASMLD %s
222+
// LINK_WASIP2_USE_WASMLD: "-cc1" {{.*}} "-o" "[[temp:[^"]*]]"
223+
// LINK_WASIP2_USE_WASMLD: wasm-ld{{.*}}" "-m" "wasm32" {{.*}} "[[temp]]" {{.*}}

‎clang/unittests/Format/FormatTest.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11475,10 +11475,10 @@ TEST_F(FormatTest, UnderstandsNewAndDelete) {
1147511475
"void new (link p);\n"
1147611476
"void delete (link p);");
1147711477

11478-
verifyFormat("p->new();\n"
11479-
"p->delete();",
11480-
"p->new ();\n"
11481-
"p->delete ();");
11478+
verifyFormat("{ p->new(); }\n"
11479+
"{ p->delete(); }",
11480+
"{ p->new (); }\n"
11481+
"{ p->delete (); }");
1148211482

1148311483
FormatStyle AfterPlacementOperator = getLLVMStyle();
1148411484
AfterPlacementOperator.SpaceBeforeParens = FormatStyle::SBPO_Custom;

‎flang/include/flang/Optimizer/Dialect/FIROps.td

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3131,6 +3131,16 @@ def fir_BoxOffsetOp : fir_Op<"box_offset", [NoMemoryEffect]> {
31313131
def fir_CUDAKernelOp : fir_Op<"cuda_kernel", [AttrSizedOperandSegments,
31323132
DeclareOpInterfaceMethods<LoopLikeOpInterface>]> {
31333133

3134+
let description = [{
3135+
Represent the CUDA Fortran kernel directive. The operation is a loop like
3136+
operation that represents the iteration range of the embedded loop nest.
3137+
3138+
When grid or block variadic operands are empty, a `*` only syntax was used
3139+
in the Fortran code.
3140+
If the `*` is mixed with values for either grid or block, these are
3141+
represented by a 0 constant value.
3142+
}];
3143+
31343144
let arguments = (ins
31353145
Variadic<I32>:$grid, // empty means `*`
31363146
Variadic<I32>:$block, // empty means `*`

‎flang/lib/Lower/Bridge.cpp

Lines changed: 32 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -2529,23 +2529,42 @@ class FirConverter : public Fortran::lower::AbstractConverter {
25292529
const std::optional<Fortran::parser::ScalarIntExpr> &stream =
25302530
std::get<3>(dir.t);
25312531

2532+
auto isOnlyStars =
2533+
[&](const std::list<Fortran::parser::CUFKernelDoConstruct::StarOrExpr>
2534+
&list) -> bool {
2535+
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr :
2536+
list) {
2537+
if (expr.v)
2538+
return false;
2539+
}
2540+
return true;
2541+
};
2542+
2543+
mlir::Value zero =
2544+
builder->createIntegerConstant(loc, builder->getI32Type(), 0);
2545+
25322546
llvm::SmallVector<mlir::Value> gridValues;
2533-
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr : grid) {
2534-
if (expr.v) {
2535-
gridValues.push_back(fir::getBase(
2536-
genExprValue(*Fortran::semantics::GetExpr(*expr.v), stmtCtx)));
2537-
} else {
2538-
// TODO: '*'
2547+
if (!isOnlyStars(grid)) {
2548+
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr :
2549+
grid) {
2550+
if (expr.v) {
2551+
gridValues.push_back(fir::getBase(
2552+
genExprValue(*Fortran::semantics::GetExpr(*expr.v), stmtCtx)));
2553+
} else {
2554+
gridValues.push_back(zero);
2555+
}
25392556
}
25402557
}
25412558
llvm::SmallVector<mlir::Value> blockValues;
2542-
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr :
2543-
block) {
2544-
if (expr.v) {
2545-
blockValues.push_back(fir::getBase(
2546-
genExprValue(*Fortran::semantics::GetExpr(*expr.v), stmtCtx)));
2547-
} else {
2548-
// TODO: '*'
2559+
if (!isOnlyStars(block)) {
2560+
for (const Fortran::parser::CUFKernelDoConstruct::StarOrExpr &expr :
2561+
block) {
2562+
if (expr.v) {
2563+
blockValues.push_back(fir::getBase(
2564+
genExprValue(*Fortran::semantics::GetExpr(*expr.v), stmtCtx)));
2565+
} else {
2566+
blockValues.push_back(zero);
2567+
}
25492568
}
25502569
}
25512570
mlir::Value streamValue;

‎flang/lib/Semantics/check-do-forall.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -438,6 +438,18 @@ class DoContext {
438438
CheckForallIndexesUsed(*assignment);
439439
CheckForImpureCall(assignment->lhs);
440440
CheckForImpureCall(assignment->rhs);
441+
442+
if (IsVariable(assignment->lhs)) {
443+
if (const Symbol * symbol{GetLastSymbol(assignment->lhs)}) {
444+
if (auto impureFinal{
445+
HasImpureFinal(*symbol, assignment->lhs.Rank())}) {
446+
context_.SayWithDecl(*symbol, parser::FindSourceLocation(stmt),
447+
"Impure procedure '%s' is referenced by finalization in a %s"_err_en_US,
448+
impureFinal->name(), LoopKindName());
449+
}
450+
}
451+
}
452+
441453
if (const auto *proc{
442454
std::get_if<evaluate::ProcedureRef>(&assignment->u)}) {
443455
CheckForImpureCall(*proc);

‎flang/test/Lower/CUDA/cuda-kernel-loop-directive.cuf

Lines changed: 16 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,20 @@ subroutine sub1()
4242
! CHECK: fir.cuda_kernel<<<%c1{{.*}}, (%c256{{.*}}, %c1{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)
4343
! CHECK: {n = 2 : i64}
4444

45-
! TODO: lowering for these cases
46-
! !$cuf kernel do(2) <<< (1,*), (256,1) >>>
47-
! !$cuf kernel do(2) <<< (*,*), (32,4) >>>
45+
!$cuf kernel do(2) <<< (1,*), (256,1) >>>
46+
do i = 1, n
47+
do j = 1, n
48+
c(i,j) = c(i,j) * d(i,j)
49+
end do
50+
end do
51+
! CHECK: fir.cuda_kernel<<<(%c1{{.*}}, %c0{{.*}}), (%c256{{.*}}, %c1{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)
52+
53+
!$cuf kernel do(2) <<< (*,*), (32,4) >>>
54+
do i = 1, n
55+
do j = 1, n
56+
c(i,j) = c(i,j) * d(i,j)
57+
end do
58+
end do
59+
60+
! CHECK: fir.cuda_kernel<<<*, (%c32{{.*}}, %c4{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)
4861
end

‎flang/test/Semantics/forall02.f90

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
1+
! RUN: %python %S/test_errors.py %s %flang_fc1
2+
3+
module m1
4+
type :: impureFinal
5+
contains
6+
final :: impureSub
7+
final :: impureSubRank1
8+
final :: impureSubRank2
9+
end type
10+
11+
contains
12+
13+
impure subroutine impureSub(x)
14+
type(impureFinal), intent(in) :: x
15+
end subroutine
16+
17+
impure subroutine impureSubRank1(x)
18+
type(impureFinal), intent(in) :: x(:)
19+
end subroutine
20+
21+
impure subroutine impureSubRank2(x)
22+
type(impureFinal), intent(in) :: x(:,:)
23+
end subroutine
24+
25+
subroutine s1()
26+
implicit none
27+
integer :: i
28+
type(impureFinal), allocatable :: ifVar, ifvar1
29+
type(impureFinal), allocatable :: ifArr1(:), ifArr2(:,:)
30+
type(impureFinal) :: if0
31+
integer a(10)
32+
allocate(ifVar)
33+
allocate(ifVar1)
34+
allocate(ifArr1(5), ifArr2(5,5))
35+
36+
! Error to invoke an IMPURE FINAL procedure in a FORALL
37+
forall (i = 1:10)
38+
!WARNING: FORALL index variable 'i' not used on left-hand side of assignment
39+
!ERROR: Impure procedure 'impuresub' is referenced by finalization in a FORALL
40+
ifvar = ifvar1
41+
end forall
42+
43+
forall (i = 1:5)
44+
!ERROR: Impure procedure 'impuresub' is referenced by finalization in a FORALL
45+
ifArr1(i) = if0
46+
end forall
47+
48+
forall (i = 1:5)
49+
!WARNING: FORALL index variable 'i' not used on left-hand side of assignment
50+
!ERROR: Impure procedure 'impuresubrank1' is referenced by finalization in a FORALL
51+
ifArr1 = if0
52+
end forall
53+
54+
forall (i = 1:5)
55+
!ERROR: Impure procedure 'impuresubrank1' is referenced by finalization in a FORALL
56+
ifArr2(i,:) = if0
57+
end forall
58+
59+
forall (i = 1:5)
60+
!WARNING: FORALL index variable 'i' not used on left-hand side of assignment
61+
!ERROR: Impure procedure 'impuresubrank2' is referenced by finalization in a FORALL
62+
ifArr2(:,:) = if0
63+
end forall
64+
end subroutine
65+
66+
end module m1
67+

‎libc/src/math/generic/CMakeLists.txt

Lines changed: 5 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -2155,16 +2155,9 @@ add_object_library(
21552155
SRCS
21562156
inv_trigf_utils.cpp
21572157
DEPENDS
2158-
.math_utils
2159-
libc.src.__support.FPUtil.fp_bits
2160-
libc.src.__support.FPUtil.fenv_impl
2161-
libc.src.__support.FPUtil.nearest_integer
2162-
libc.src.__support.FPUtil.nearest_integer_operations
2158+
libc.src.__support.FPUtil.multiply_add
21632159
libc.src.__support.FPUtil.polyeval
21642160
libc.src.__support.common
2165-
libc.include.errno
2166-
libc.src.errno.errno
2167-
libc.include.math
21682161
)
21692162

21702163
add_entrypoint_object(
@@ -2211,8 +2204,11 @@ add_entrypoint_object(
22112204
../atanf.h
22122205
DEPENDS
22132206
.inv_trigf_utils
2214-
.math_utils
2207+
libc.src.__support.FPUtil.except_value_utils
22152208
libc.src.__support.FPUtil.fp_bits
2209+
libc.src.__support.FPUtil.multiply_add
2210+
libc.src.__support.FPUtil.nearest_integer
2211+
libc.src.__support.FPUtil.polyeval
22162212
libc.src.__support.FPUtil.rounding_mode
22172213
libc.src.__support.macros.optimization
22182214
COMPILE_OPTIONS

0 commit comments

Comments
 (0)