Skip to content

Commit b88886b

Browse files
hekotahekota
authored
Merge branch 'main' into cbuffer-layout-structs
2 parents 71ddb5a + ffe5cdd commit b88886b

File tree

87 files changed

+12159
-10440
lines changed

Some content is hidden

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

87 files changed

+12159
-10440
lines changed

‎clang/lib/Analysis/UnsafeBufferUsage.cpp

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -453,8 +453,11 @@ AST_MATCHER(ArraySubscriptExpr, isSafeArraySubscript) {
453453
return false;
454454
}
455455

456-
if (const auto *IdxLit = dyn_cast<IntegerLiteral>(Node.getIdx())) {
457-
const APInt ArrIdx = IdxLit->getValue();
456+
Expr::EvalResult EVResult;
457+
if (Node.getIdx()->EvaluateAsInt(EVResult, Finder->getASTContext())) {
458+
llvm::APSInt ArrIdx = EVResult.Val.getInt();
459+
// FIXME: ArrIdx.isNegative() we could immediately emit an error as that's a
460+
// bug
458461
if (ArrIdx.isNonNegative() && ArrIdx.getLimitedValue() < limit)
459462
return true;
460463
}

‎clang/test/SemaCXX/warn-unsafe-buffer-usage-array.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -92,3 +92,35 @@ char access_strings() {
9292
c = array_string[5];
9393
return c;
9494
}
95+
96+
struct T {
97+
int array[10];
98+
};
99+
100+
const int index = 1;
101+
102+
constexpr int get_const(int x) {
103+
if(x < 3)
104+
return ++x;
105+
else
106+
return x + 5;
107+
};
108+
109+
void array_indexed_const_expr(unsigned idx) {
110+
// expected-note@+2 {{change type of 'arr' to 'std::array' to label it for hardening}}
111+
// expected-warning@+1{{'arr' is an unsafe buffer that does not perform bounds checks}}
112+
int arr[10];
113+
arr[sizeof(int)] = 5;
114+
115+
int array[sizeof(T)];
116+
array[sizeof(int)] = 5;
117+
array[sizeof(T) -1 ] = 3;
118+
119+
int k = arr[6 & 5];
120+
k = arr[2 << index];
121+
k = arr[8 << index]; // expected-note {{used in buffer access here}}
122+
k = arr[16 >> 1];
123+
k = arr[get_const(index)];
124+
k = arr[get_const(5)]; // expected-note {{used in buffer access here}}
125+
k = arr[get_const(4)];
126+
}

‎flang/lib/Optimizer/Transforms/CUFOpConversion.cpp

Lines changed: 44 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -366,6 +366,23 @@ struct CUFAllocOpConversion : public mlir::OpRewritePattern<cuf::AllocOp> {
366366
const fir::LLVMTypeConverter *typeConverter;
367367
};
368368

369+
static mlir::Value genGetDeviceAddress(mlir::PatternRewriter &rewriter,
370+
mlir::ModuleOp mod, mlir::Location loc,
371+
mlir::Value inputArg) {
372+
fir::FirOpBuilder builder(rewriter, mod);
373+
mlir::func::FuncOp callee =
374+
fir::runtime::getRuntimeFunc<mkRTKey(CUFGetDeviceAddress)>(loc, builder);
375+
auto fTy = callee.getFunctionType();
376+
mlir::Value conv = createConvertOp(rewriter, loc, fTy.getInput(0), inputArg);
377+
mlir::Value sourceFile = fir::factory::locationToFilename(builder, loc);
378+
mlir::Value sourceLine =
379+
fir::factory::locationToLineNo(builder, loc, fTy.getInput(2));
380+
llvm::SmallVector<mlir::Value> args{fir::runtime::createArguments(
381+
builder, loc, fTy, conv, sourceFile, sourceLine)};
382+
auto call = rewriter.create<fir::CallOp>(loc, callee, args);
383+
return createConvertOp(rewriter, loc, inputArg.getType(), call->getResult(0));
384+
}
385+
369386
struct DeclareOpConversion : public mlir::OpRewritePattern<fir::DeclareOp> {
370387
using OpRewritePattern::OpRewritePattern;
371388

@@ -382,26 +399,10 @@ struct DeclareOpConversion : public mlir::OpRewritePattern<fir::DeclareOp> {
382399
if (cuf::isRegisteredDeviceGlobal(global)) {
383400
rewriter.setInsertionPointAfter(addrOfOp);
384401
auto mod = op->getParentOfType<mlir::ModuleOp>();
385-
fir::FirOpBuilder builder(rewriter, mod);
386-
mlir::Location loc = op.getLoc();
387-
mlir::func::FuncOp callee =
388-
fir::runtime::getRuntimeFunc<mkRTKey(CUFGetDeviceAddress)>(
389-
loc, builder);
390-
auto fTy = callee.getFunctionType();
391-
mlir::Type toTy = fTy.getInput(0);
392-
mlir::Value inputArg =
393-
createConvertOp(rewriter, loc, toTy, addrOfOp.getResult());
394-
mlir::Value sourceFile =
395-
fir::factory::locationToFilename(builder, loc);
396-
mlir::Value sourceLine =
397-
fir::factory::locationToLineNo(builder, loc, fTy.getInput(2));
398-
llvm::SmallVector<mlir::Value> args{fir::runtime::createArguments(
399-
builder, loc, fTy, inputArg, sourceFile, sourceLine)};
400-
auto call = rewriter.create<fir::CallOp>(loc, callee, args);
401-
mlir::Value cast = createConvertOp(
402-
rewriter, loc, op.getMemref().getType(), call->getResult(0));
402+
mlir::Value devAddr = genGetDeviceAddress(rewriter, mod, op.getLoc(),
403+
addrOfOp.getResult());
403404
rewriter.startOpModification(op);
404-
op.getMemrefMutable().assign(cast);
405+
op.getMemrefMutable().assign(devAddr);
405406
rewriter.finalizeOpModification(op);
406407
return success();
407408
}
@@ -771,10 +772,32 @@ struct CUFLaunchOpConversion
771772
loc, clusterDimsAttr.getZ().getInt());
772773
}
773774
}
775+
llvm::SmallVector<mlir::Value> args;
776+
auto mod = op->getParentOfType<mlir::ModuleOp>();
777+
for (mlir::Value arg : op.getArgs()) {
778+
// If the argument is a global descriptor, make sure we pass the device
779+
// copy of this descriptor and not the host one.
780+
if (mlir::isa<fir::BaseBoxType>(fir::unwrapRefType(arg.getType()))) {
781+
if (auto declareOp =
782+
mlir::dyn_cast_or_null<fir::DeclareOp>(arg.getDefiningOp())) {
783+
if (auto addrOfOp = mlir::dyn_cast_or_null<fir::AddrOfOp>(
784+
declareOp.getMemref().getDefiningOp())) {
785+
if (auto global = symTab.lookup<fir::GlobalOp>(
786+
addrOfOp.getSymbol().getRootReference().getValue())) {
787+
if (cuf::isRegisteredDeviceGlobal(global)) {
788+
arg = genGetDeviceAddress(rewriter, mod, op.getLoc(),
789+
declareOp.getResult());
790+
}
791+
}
792+
}
793+
}
794+
}
795+
args.push_back(arg);
796+
}
797+
774798
auto gpuLaunchOp = rewriter.create<mlir::gpu::LaunchFuncOp>(
775799
loc, kernelName, mlir::gpu::KernelDim3{gridSizeX, gridSizeY, gridSizeZ},
776-
mlir::gpu::KernelDim3{blockSizeX, blockSizeY, blockSizeZ}, zero,
777-
op.getArgs());
800+
mlir::gpu::KernelDim3{blockSizeX, blockSizeY, blockSizeZ}, zero, args);
778801
if (clusterDimX && clusterDimY && clusterDimZ) {
779802
gpuLaunchOp.getClusterSizeXMutable().assign(clusterDimX);
780803
gpuLaunchOp.getClusterSizeYMutable().assign(clusterDimY);

‎flang/lib/Semantics/resolve-names.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4015,7 +4015,9 @@ bool SubprogramVisitor::Pre(const parser::PrefixSpec::Attributes &attrs) {
40154015
*attrs == common::CUDASubprogramAttrs::Device) {
40164016
const Scope &scope{currScope()};
40174017
const Scope *mod{FindModuleContaining(scope)};
4018-
if (mod && mod->GetName().value() == "cudadevice") {
4018+
if (mod &&
4019+
(mod->GetName().value() == "cudadevice" ||
4020+
mod->GetName().value() == "__cuda_device")) {
40194021
return false;
40204022
}
40214023
// Implicitly USE the cudadevice module by copying its symbols in the

‎flang/module/__cuda_device.f90

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
!===-- module/__cuda_device.f90 --------------------------------------------===!
2+
!
3+
! Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
! See https://llvm.org/LICENSE.txt for license information.
5+
! SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
!
7+
!===------------------------------------------------------------------------===!
8+
9+
! This module contains CUDA Fortran interfaces used in cudadevice.f90.
10+
11+
module __cuda_device
12+
implicit none
13+
14+
! Set PRIVATE by default to explicitly only export what is meant
15+
! to be exported by this MODULE.
16+
17+
interface
18+
attributes(device) function __fadd_rd(x, y) bind(c, name='__nv_fadd_rd')
19+
real, intent(in), value :: x, y
20+
real :: __fadd_rd
21+
end function
22+
end interface
23+
public :: __fadd_rd
24+
25+
interface
26+
attributes(device) function __fadd_ru(x, y) bind(c, name='__nv_fadd_ru')
27+
real, intent(in), value :: x, y
28+
real :: __fadd_ru
29+
end function
30+
end interface
31+
public :: __fadd_ru
32+
end module

‎flang/module/cudadevice.f90

Lines changed: 1 addition & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
! CUDA Fortran procedures available in device subprogram
1010

1111
module cudadevice
12+
use __cuda_device, only: __fadd_rd, __fadd_ru
1213
implicit none
1314

1415
! Set PRIVATE by default to explicitly only export what is meant
@@ -71,20 +72,4 @@ attributes(device) subroutine threadfence_system()
7172
end interface
7273
public :: threadfence_system
7374

74-
interface
75-
attributes(device) function __fadd_rd(x, y) bind(c, name='__nv_fadd_rd')
76-
real, intent(in) :: x, y
77-
real :: __fadd_rd
78-
end function
79-
end interface
80-
public :: __fadd_rd
81-
82-
interface
83-
attributes(device) function __fadd_ru(x, y) bind(c, name='__nv_fadd_ru')
84-
real, intent(in) :: x, y
85-
real :: __fadd_ru
86-
end function
87-
end interface
88-
public :: __fadd_ru
89-
9075
end module

‎flang/test/Fir/CUDA/cuda-launch.fir

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -62,3 +62,45 @@ module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_e
6262
// CHECK-LABEL: func.func @_QMmod1Phost_sub()
6363
// CHECK: gpu.launch_func @cuda_device_mod::@_QMmod1Psub1 clusters in (%c2{{.*}}, %c2{{.*}}, %c1{{.*}})
6464

65+
// -----
66+
67+
module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
68+
gpu.module @cuda_device_mod {
69+
gpu.func @_QMdevptrPtest(%arg0: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>) kernel {
70+
gpu.return
71+
}
72+
}
73+
fir.global @_QMdevptrEdev_ptr {data_attr = #cuf.cuda<device>} : !fir.box<!fir.ptr<!fir.array<?xf32>>> {
74+
%c0 = arith.constant 0 : index
75+
%0 = fir.zero_bits !fir.ptr<!fir.array<?xf32>>
76+
%1 = fir.shape %c0 : (index) -> !fir.shape<1>
77+
%2 = fir.embox %0(%1) {allocator_idx = 2 : i32} : (!fir.ptr<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.ptr<!fir.array<?xf32>>>
78+
fir.has_value %2 : !fir.box<!fir.ptr<!fir.array<?xf32>>>
79+
}
80+
func.func @_QMdevptrPtest(%arg0: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>> {cuf.data_attr = #cuf.cuda<device>, fir.bindc_name = "dp"}) attributes {cuf.proc_attr = #cuf.cuda_proc<global>} {
81+
return
82+
}
83+
func.func @_QQmain() {
84+
%c1_i32 = arith.constant 1 : i32
85+
%c4 = arith.constant 4 : index
86+
%0 = cuf.alloc !fir.array<4xf32> {bindc_name = "a_dev", data_attr = #cuf.cuda<device>, uniq_name = "_QFEa_dev"} -> !fir.ref<!fir.array<4xf32>>
87+
%1 = fir.shape %c4 : (index) -> !fir.shape<1>
88+
%2 = fir.declare %0(%1) {data_attr = #cuf.cuda<device>, fortran_attrs = #fir.var_attrs<target>, uniq_name = "_QFEa_dev"} : (!fir.ref<!fir.array<4xf32>>, !fir.shape<1>) -> !fir.ref<!fir.array<4xf32>>
89+
%3 = fir.address_of(@_QMdevptrEdev_ptr) : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>
90+
%4 = fir.declare %3 {data_attr = #cuf.cuda<device>, fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QMdevptrEdev_ptr"} : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>
91+
%5 = fir.embox %2(%1) : (!fir.ref<!fir.array<4xf32>>, !fir.shape<1>) -> !fir.box<!fir.ptr<!fir.array<?xf32>>>
92+
fir.store %5 to %4 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>
93+
cuf.sync_descriptor @_QMdevptrEdev_ptr
94+
cuf.kernel_launch @_QMdevptrPtest<<<%c1_i32, %c1_i32, %c1_i32, %c1_i32, %c1_i32, %c1_i32>>>(%4) : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>)
95+
cuf.free %2 : !fir.ref<!fir.array<4xf32>> {data_attr = #cuf.cuda<device>}
96+
return
97+
}
98+
}
99+
100+
// CHECK-LABEL: func.func @_QQmain()
101+
// CHECK: %[[ADDROF:.*]] = fir.address_of(@_QMdevptrEdev_ptr) : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>
102+
// CHECK: %[[DECL:.*]] = fir.declare %[[ADDROF]] {data_attr = #cuf.cuda<device>, fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QMdevptrEdev_ptr"} : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>
103+
// CHECK: %[[CONV_DECL:.*]] = fir.convert %[[DECL]] : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>) -> !fir.llvm_ptr<i8>
104+
// CHECK: %[[DEVADDR:.*]] = fir.call @_FortranACUFGetDeviceAddress(%[[CONV_DECL]], %{{.*}}, %{{.*}}) : (!fir.llvm_ptr<i8>, !fir.ref<i8>, i32) -> !fir.llvm_ptr<i8>
105+
// CHECK: %[[CONV_DEVADDR:.*]] = fir.convert %[[DEVADDR]] : (!fir.llvm_ptr<i8>) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>
106+
// CHECK: gpu.launch_func @cuda_device_mod::@_QMdevptrPtest blocks in (%{{.*}}, %{{.*}}, %{{.*}}) threads in (%{{.*}}, %{{.*}}, %{{.*}}) dynamic_shared_memory_size %{{.*}} args(%[[CONV_DEVADDR]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>)

‎flang/tools/f18/CMakeLists.txt

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@ set(MODULES_WITHOUT_IMPLEMENTATION
2121
"__ppc_intrinsics"
2222
"mma"
2323
"__cuda_builtins"
24+
"__cuda_device"
2425
"cudadevice"
2526
"ieee_arithmetic"
2627
"ieee_exceptions"
@@ -67,9 +68,12 @@ if (NOT CMAKE_CROSSCOMPILING)
6768
elseif(${filename} STREQUAL "__ppc_intrinsics" OR
6869
${filename} STREQUAL "mma")
6970
set(depends ${FLANG_INTRINSIC_MODULES_DIR}/__ppc_types.mod)
70-
elseif(${filename} STREQUAL "cudadevice")
71+
elseif(${filename} STREQUAL "__cuda_device")
7172
set(opts -fc1 -xcuda)
7273
set(depends ${FLANG_INTRINSIC_MODULES_DIR}/__cuda_builtins.mod)
74+
elseif(${filename} STREQUAL "cudadevice")
75+
set(opts -fc1 -xcuda)
76+
set(depends ${FLANG_INTRINSIC_MODULES_DIR}/__cuda_device.mod)
7377
else()
7478
set(depends ${FLANG_INTRINSIC_MODULES_DIR}/__fortran_builtins.mod)
7579
if(${filename} STREQUAL "iso_fortran_env")

‎lldb/packages/Python/lldbsuite/test/lldbtest.py

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1379,6 +1379,17 @@ def isAArch64Windows(self):
13791379
return arch in ["aarch64", "arm64", "arm64e"]
13801380
return False
13811381

1382+
def isLoongArch(self):
1383+
"""Returns true if the architecture is LoongArch."""
1384+
arch = self.getArchitecture().lower()
1385+
return arch in ["loongarch64", "loongarch32"]
1386+
1387+
def isLoongArchLSX(self):
1388+
return self.isLoongArch() and "lsx" in self.getCPUInfo()
1389+
1390+
def isLoongArchLASX(self):
1391+
return self.isLoongArch() and "lasx" in self.getCPUInfo()
1392+
13821393
def getArchitecture(self):
13831394
"""Returns the architecture in effect the test suite is running with."""
13841395
return lldbplatformutil.getArchitecture()

‎lldb/packages/Python/lldbsuite/test/make/Makefile.rules

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -210,6 +210,10 @@ else
210210
ifeq "$(findstring mips,$(ARCH))" "mips"
211211
override ARCHFLAG := -
212212
endif
213+
ifeq "$(findstring loongarch,$(ARCH))" "loongarch"
214+
override ARCH :=
215+
override ARCHFLAG :=
216+
endif
213217

214218
ifeq "$(SPLIT_DEBUG_SYMBOLS)" "YES"
215219
DSYM = $(EXE).debug

0 commit comments

Comments
 (0)