Skip to content

Commit e8c2dcb

Browse files
SC llvm teamSC llvm team
SC llvm team
authored and
SC llvm team
committed
Merge upstream LLVM into amd-gfx12
2 parents 0e20ee0 + ae3bba4 commit e8c2dcb

File tree

24 files changed

+644
-144
lines changed

24 files changed

+644
-144
lines changed

‎clang/lib/CodeGen/CGCoroutine.cpp‎

Lines changed: 9 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -707,11 +707,15 @@ struct GetReturnObjectManager {
707707
Builder.CreateStore(Builder.getFalse(), GroActiveFlag);
708708

709709
GroEmission = CGF.EmitAutoVarAlloca(*GroVarDecl);
710-
auto *GroAlloca = dyn_cast_or_null<llvm::AllocaInst>(
711-
GroEmission.getOriginalAllocatedAddress().getPointer());
712-
assert(GroAlloca && "expected alloca to be emitted");
713-
GroAlloca->setMetadata(llvm::LLVMContext::MD_coro_outside_frame,
714-
llvm::MDNode::get(CGF.CGM.getLLVMContext(), {}));
710+
711+
if (!GroVarDecl->isNRVOVariable()) {
712+
// NRVO variables don't have allocas and won't have the same issue.
713+
auto *GroAlloca = dyn_cast_or_null<llvm::AllocaInst>(
714+
GroEmission.getOriginalAllocatedAddress().getPointer());
715+
assert(GroAlloca && "expected alloca to be emitted");
716+
GroAlloca->setMetadata(llvm::LLVMContext::MD_coro_outside_frame,
717+
llvm::MDNode::get(CGF.CGM.getLLVMContext(), {}));
718+
}
715719

716720
// Remember the top of EHStack before emitting the cleanup.
717721
auto old_top = CGF.EHStack.stable_begin();

‎clang/lib/Headers/CMakeLists.txt‎

Lines changed: 19 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -347,6 +347,10 @@ set(cuda_wrapper_bits_files
347347
cuda_wrappers/bits/basic_string.tcc
348348
)
349349

350+
set(cuda_wrapper_utility_files
351+
cuda_wrappers/__utility/declval.h
352+
)
353+
350354
set(ppc_wrapper_files
351355
ppc_wrappers/mmintrin.h
352356
ppc_wrappers/xmmintrin.h
@@ -443,8 +447,9 @@ endfunction(clang_generate_header)
443447

444448
# Copy header files from the source directory to the build directory
445449
foreach( f ${files} ${cuda_wrapper_files} ${cuda_wrapper_bits_files}
446-
${ppc_wrapper_files} ${openmp_wrapper_files} ${zos_wrapper_files} ${hlsl_files}
447-
${llvm_libc_wrapper_files} ${llvm_offload_wrapper_files})
450+
${cuda_wrapper_utility_files} ${ppc_wrapper_files} ${openmp_wrapper_files}
451+
${zos_wrapper_files} ${hlsl_files} ${llvm_libc_wrapper_files}
452+
${llvm_offload_wrapper_files})
448453
copy_header_to_output_dir(${CMAKE_CURRENT_SOURCE_DIR} ${f})
449454
endforeach( f )
450455

@@ -553,7 +558,7 @@ add_header_target("arm-common-resource-headers" "${arm_common_files};${arm_commo
553558
# Architecture/platform specific targets
554559
add_header_target("arm-resource-headers" "${arm_only_files};${arm_only_generated_files}")
555560
add_header_target("aarch64-resource-headers" "${aarch64_only_files};${aarch64_only_generated_files}")
556-
add_header_target("cuda-resource-headers" "${cuda_files};${cuda_wrapper_files};${cuda_wrapper_bits_files}")
561+
add_header_target("cuda-resource-headers" "${cuda_files};${cuda_wrapper_files};${cuda_wrapper_bits_files};${cuda_wrapper_utility_files}")
557562
add_header_target("hexagon-resource-headers" "${hexagon_files}")
558563
add_header_target("hip-resource-headers" "${hip_files}")
559564
add_header_target("loongarch-resource-headers" "${loongarch_files}")
@@ -600,6 +605,11 @@ install(
600605
DESTINATION ${header_install_dir}/cuda_wrappers/bits
601606
COMPONENT clang-resource-headers)
602607

608+
install(
609+
FILES ${cuda_wrapper_utility_files}
610+
DESTINATION ${header_install_dir}/cuda_wrappers/__utility
611+
COMPONENT clang-resource-headers)
612+
603613
install(
604614
FILES ${ppc_wrapper_files}
605615
DESTINATION ${header_install_dir}/ppc_wrappers
@@ -663,6 +673,12 @@ install(
663673
EXCLUDE_FROM_ALL
664674
COMPONENT cuda-resource-headers)
665675

676+
install(
677+
FILES ${cuda_wrapper_utility_files}
678+
DESTINATION ${header_install_dir}/cuda_wrappers/__utility
679+
EXCLUDE_FROM_ALL
680+
COMPONENT cuda-resource-headers)
681+
666682
install(
667683
FILES ${cuda_files}
668684
DESTINATION ${header_install_dir}

‎clang/lib/Headers/cuda_wrappers/__utility/declval.h‎

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
#ifndef __CUDA_WRAPPERS_UTILITY_DECLVAL_H__
2+
#define __CUDA_WRAPPERS_UTILITY_DECLVAL_H__
3+
4+
#include_next <__utility/declval.h>
5+
6+
// The stuff below is the exact copy of the <__utility/declval.h>,
7+
// but with __device__ attribute applied to the functions, so it works on a GPU.
8+
9+
_LIBCPP_BEGIN_NAMESPACE_STD
10+
11+
// Suppress deprecation notice for volatile-qualified return type resulting
12+
// from volatile-qualified types _Tp.
13+
_LIBCPP_SUPPRESS_DEPRECATED_PUSH
14+
template <class _Tp> __attribute__((device)) _Tp &&__declval(int);
15+
template <class _Tp> __attribute__((device)) _Tp __declval(long);
16+
_LIBCPP_SUPPRESS_DEPRECATED_POP
17+
18+
template <class _Tp>
19+
__attribute__((device)) _LIBCPP_HIDE_FROM_ABI decltype(std::__declval<_Tp>(0))
20+
declval() _NOEXCEPT {
21+
static_assert(!__is_same(_Tp, _Tp),
22+
"std::declval can only be used in an unevaluated context. "
23+
"It's likely that your current usage is trying to extract a "
24+
"value from the function.");
25+
}
26+
27+
_LIBCPP_END_NAMESPACE_STD
28+
#endif // __CUDA_WRAPPERS_UTILITY_DECLVAL_H__

‎clang/lib/Sema/SemaHLSL.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1150,7 +1150,7 @@ bool SemaHLSL::handleRootSignatureElements(
11501150
if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
11511151
ReportError(Loc, 0, 16);
11521152
if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1153-
ReportFloatError(Loc, -16.f, 15.99);
1153+
ReportFloatError(Loc, -16.f, 15.99f);
11541154
} else if (const auto *Clause =
11551155
std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
11561156
&Elem)) {

‎clang/lib/StaticAnalyzer/Checkers/WebKit/ASTUtils.cpp‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -177,7 +177,10 @@ bool tryToFindPtrOrigin(
177177
E = unaryOp->getSubExpr();
178178
continue;
179179
}
180-
180+
if (auto *BoxedExpr = dyn_cast<ObjCBoxedExpr>(E)) {
181+
E = BoxedExpr->getSubExpr();
182+
continue;
183+
}
181184
break;
182185
}
183186
// Some other expression.

‎clang/test/Analysis/Checkers/WebKit/unretained-call-args.mm‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -329,13 +329,17 @@ void foo() {
329329
}
330330
}
331331

332+
#define YES 1
333+
332334
namespace call_with_cf_constant {
333335
void bar(const NSArray *);
334336
void baz(const NSDictionary *);
337+
void boo(NSNumber *);
335338
void foo() {
336339
CFArrayCreateMutable(kCFAllocatorDefault, 10);
337340
bar(@[@"hello"]);
338341
baz(@{@"hello": @3});
342+
boo(@YES);
339343
}
340344
}
341345

‎clang/test/CodeGenCoroutines/coro-gro.cpp‎

Lines changed: 28 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -106,4 +106,31 @@ invoker g() {
106106
// CHECK: call void @_ZN7invoker15invoker_promise17get_return_objectEv({{.*}} %[[AggRes]]
107107
co_return;
108108
}
109-
// CHECK: ![[OutFrameMetadata]] = !{}
109+
110+
namespace gh148953 {
111+
112+
struct Task {
113+
struct promise_type {
114+
Task get_return_object();
115+
std::suspend_always initial_suspend() { return {}; }
116+
std::suspend_always final_suspend() noexcept { return {}; }
117+
void return_void() {}
118+
void unhandled_exception() {}
119+
};
120+
Task() {}
121+
// Different from `invoker`, this Task is copy constructible.
122+
Task(const Task&) {};
123+
};
124+
125+
// NRVO on const qualified return type should work.
126+
// CHECK: define{{.*}} void @_ZN8gh1489537exampleEv({{.*}} sret(%"struct.gh148953::Task") align 1 %[[NrvoRes:.+]])
127+
const Task example() {
128+
// CHECK: %[[ResultPtr:.+]] = alloca ptr
129+
// CHECK: store ptr %[[NrvoRes]], ptr %[[ResultPtr]]
130+
// CHECK: coro.init:
131+
// CHECK: call void @_ZN8gh1489534Task12promise_type17get_return_objectEv({{.*}} %[[NrvoRes:.+]], {{.*}})
132+
co_return;
133+
}
134+
135+
} // namespace gh148953
136+
// CHECK: ![[OutFrameMetadata]] = !{}

‎flang/lib/Lower/ConvertVariable.cpp‎

Lines changed: 94 additions & 73 deletions
Original file line numberDiff line numberDiff line change
@@ -771,79 +771,9 @@ static mlir::Value createNewLocal(Fortran::lower::AbstractConverter &converter,
771771
return builder.create<cuf::SharedMemoryOp>(loc, ty, nm, symNm, lenParams,
772772
indices);
773773

774-
if (!cuf::isCUDADeviceContext(builder.getRegion())) {
775-
mlir::Value alloc = builder.create<cuf::AllocOp>(
776-
loc, ty, nm, symNm, dataAttr, lenParams, indices);
777-
if (const auto *details{
778-
ultimateSymbol
779-
.detailsIf<Fortran::semantics::ObjectEntityDetails>()}) {
780-
const Fortran::semantics::DeclTypeSpec *type{details->type()};
781-
const Fortran::semantics::DerivedTypeSpec *derived{
782-
type ? type->AsDerived() : nullptr};
783-
if (derived) {
784-
Fortran::semantics::UltimateComponentIterator components{*derived};
785-
auto recTy = mlir::dyn_cast<fir::RecordType>(ty);
786-
787-
llvm::SmallVector<mlir::Value> coordinates;
788-
for (const auto &sym : components) {
789-
if (Fortran::semantics::IsDeviceAllocatable(sym)) {
790-
unsigned fieldIdx = recTy.getFieldIndex(sym.name().ToString());
791-
mlir::Type fieldTy;
792-
std::vector<mlir::Value> coordinates;
793-
794-
if (fieldIdx != std::numeric_limits<unsigned>::max()) {
795-
// Field found in the base record type.
796-
auto fieldName = recTy.getTypeList()[fieldIdx].first;
797-
fieldTy = recTy.getTypeList()[fieldIdx].second;
798-
mlir::Value fieldIndex = builder.create<fir::FieldIndexOp>(
799-
loc, fir::FieldType::get(fieldTy.getContext()), fieldName,
800-
recTy,
801-
/*typeParams=*/mlir::ValueRange{});
802-
coordinates.push_back(fieldIndex);
803-
} else {
804-
// Field not found in base record type, search in potential
805-
// record type components.
806-
for (auto component : recTy.getTypeList()) {
807-
if (auto childRecTy =
808-
mlir::dyn_cast<fir::RecordType>(component.second)) {
809-
fieldIdx = childRecTy.getFieldIndex(sym.name().ToString());
810-
if (fieldIdx != std::numeric_limits<unsigned>::max()) {
811-
mlir::Value parentFieldIndex =
812-
builder.create<fir::FieldIndexOp>(
813-
loc, fir::FieldType::get(childRecTy.getContext()),
814-
component.first, recTy,
815-
/*typeParams=*/mlir::ValueRange{});
816-
coordinates.push_back(parentFieldIndex);
817-
auto fieldName = childRecTy.getTypeList()[fieldIdx].first;
818-
fieldTy = childRecTy.getTypeList()[fieldIdx].second;
819-
mlir::Value childFieldIndex =
820-
builder.create<fir::FieldIndexOp>(
821-
loc, fir::FieldType::get(fieldTy.getContext()),
822-
fieldName, childRecTy,
823-
/*typeParams=*/mlir::ValueRange{});
824-
coordinates.push_back(childFieldIndex);
825-
break;
826-
}
827-
}
828-
}
829-
}
830-
831-
if (coordinates.empty())
832-
TODO(loc, "device resident component in complex derived-type "
833-
"hierarchy");
834-
835-
mlir::Value comp = builder.create<fir::CoordinateOp>(
836-
loc, builder.getRefType(fieldTy), alloc, coordinates);
837-
cuf::DataAttributeAttr dataAttr =
838-
Fortran::lower::translateSymbolCUFDataAttribute(
839-
builder.getContext(), sym);
840-
builder.create<cuf::SetAllocatorIndexOp>(loc, comp, dataAttr);
841-
}
842-
}
843-
}
844-
}
845-
return alloc;
846-
}
774+
if (!cuf::isCUDADeviceContext(builder.getRegion()))
775+
return builder.create<cuf::AllocOp>(loc, ty, nm, symNm, dataAttr,
776+
lenParams, indices);
847777
}
848778

849779
// Let the builder do all the heavy lifting.
@@ -857,6 +787,91 @@ static mlir::Value createNewLocal(Fortran::lower::AbstractConverter &converter,
857787
return res;
858788
}
859789

790+
/// Device allocatable components in a derived-type don't have the correct
791+
/// allocator index in their descriptor when they are created. After
792+
/// initialization, cuf.set_allocator_idx operations are inserted to set the
793+
/// correct allocator index for each device component.
794+
static void
795+
initializeDeviceComponentAllocator(Fortran::lower::AbstractConverter &converter,
796+
const Fortran::semantics::Symbol &symbol,
797+
Fortran::lower::SymMap &symMap) {
798+
if (const auto *details{
799+
symbol.GetUltimate()
800+
.detailsIf<Fortran::semantics::ObjectEntityDetails>()}) {
801+
const Fortran::semantics::DeclTypeSpec *type{details->type()};
802+
const Fortran::semantics::DerivedTypeSpec *derived{type ? type->AsDerived()
803+
: nullptr};
804+
if (derived) {
805+
fir::FirOpBuilder &builder = converter.getFirOpBuilder();
806+
mlir::Location loc = converter.getCurrentLocation();
807+
808+
fir::ExtendedValue exv =
809+
converter.getSymbolExtendedValue(symbol.GetUltimate(), &symMap);
810+
auto recTy = mlir::dyn_cast<fir::RecordType>(
811+
fir::unwrapRefType(fir::getBase(exv).getType()));
812+
assert(recTy && "expected fir::RecordType");
813+
814+
llvm::SmallVector<mlir::Value> coordinates;
815+
Fortran::semantics::UltimateComponentIterator components{*derived};
816+
for (const auto &sym : components) {
817+
if (Fortran::semantics::IsDeviceAllocatable(sym)) {
818+
unsigned fieldIdx = recTy.getFieldIndex(sym.name().ToString());
819+
mlir::Type fieldTy;
820+
std::vector<mlir::Value> coordinates;
821+
822+
if (fieldIdx != std::numeric_limits<unsigned>::max()) {
823+
// Field found in the base record type.
824+
auto fieldName = recTy.getTypeList()[fieldIdx].first;
825+
fieldTy = recTy.getTypeList()[fieldIdx].second;
826+
mlir::Value fieldIndex = builder.create<fir::FieldIndexOp>(
827+
loc, fir::FieldType::get(fieldTy.getContext()), fieldName,
828+
recTy,
829+
/*typeParams=*/mlir::ValueRange{});
830+
coordinates.push_back(fieldIndex);
831+
} else {
832+
// Field not found in base record type, search in potential
833+
// record type components.
834+
for (auto component : recTy.getTypeList()) {
835+
if (auto childRecTy =
836+
mlir::dyn_cast<fir::RecordType>(component.second)) {
837+
fieldIdx = childRecTy.getFieldIndex(sym.name().ToString());
838+
if (fieldIdx != std::numeric_limits<unsigned>::max()) {
839+
mlir::Value parentFieldIndex =
840+
builder.create<fir::FieldIndexOp>(
841+
loc, fir::FieldType::get(childRecTy.getContext()),
842+
component.first, recTy,
843+
/*typeParams=*/mlir::ValueRange{});
844+
coordinates.push_back(parentFieldIndex);
845+
auto fieldName = childRecTy.getTypeList()[fieldIdx].first;
846+
fieldTy = childRecTy.getTypeList()[fieldIdx].second;
847+
mlir::Value childFieldIndex =
848+
builder.create<fir::FieldIndexOp>(
849+
loc, fir::FieldType::get(fieldTy.getContext()),
850+
fieldName, childRecTy,
851+
/*typeParams=*/mlir::ValueRange{});
852+
coordinates.push_back(childFieldIndex);
853+
break;
854+
}
855+
}
856+
}
857+
}
858+
859+
if (coordinates.empty())
860+
TODO(loc, "device resident component in complex derived-type "
861+
"hierarchy");
862+
863+
mlir::Value comp = builder.create<fir::CoordinateOp>(
864+
loc, builder.getRefType(fieldTy), fir::getBase(exv), coordinates);
865+
cuf::DataAttributeAttr dataAttr =
866+
Fortran::lower::translateSymbolCUFDataAttribute(
867+
builder.getContext(), sym);
868+
builder.create<cuf::SetAllocatorIndexOp>(loc, comp, dataAttr);
869+
}
870+
}
871+
}
872+
}
873+
}
874+
860875
/// Must \p var be default initialized at runtime when entering its scope.
861876
static bool
862877
mustBeDefaultInitializedAtRuntime(const Fortran::lower::pft::Variable &var) {
@@ -1179,6 +1194,9 @@ static void instantiateLocal(Fortran::lower::AbstractConverter &converter,
11791194
if (mustBeDefaultInitializedAtRuntime(var))
11801195
Fortran::lower::defaultInitializeAtRuntime(converter, var.getSymbol(),
11811196
symMap);
1197+
if (converter.getFoldingContext().languageFeatures().IsEnabled(
1198+
Fortran::common::LanguageFeature::CUDA))
1199+
initializeDeviceComponentAllocator(converter, var.getSymbol(), symMap);
11821200
auto *builder = &converter.getFirOpBuilder();
11831201
if (needCUDAAlloc(var.getSymbol()) &&
11841202
!cuf::isCUDADeviceContext(builder->getRegion())) {
@@ -1437,6 +1455,9 @@ static void instantiateAlias(Fortran::lower::AbstractConverter &converter,
14371455
if (mustBeDefaultInitializedAtRuntime(var))
14381456
Fortran::lower::defaultInitializeAtRuntime(converter, var.getSymbol(),
14391457
symMap);
1458+
if (converter.getFoldingContext().languageFeatures().IsEnabled(
1459+
Fortran::common::LanguageFeature::CUDA))
1460+
initializeDeviceComponentAllocator(converter, var.getSymbol(), symMap);
14401461
}
14411462

14421463
//===--------------------------------------------------------------===//

‎flang/test/Lower/CUDA/cuda-set-allocator.cuf‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -12,10 +12,13 @@ contains
1212
end subroutine
1313

1414
! CHECK-LABEL: func.func @_QMm1Psub1()
15-
! CHECK: %[[DT:.*]] = cuf.alloc !fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}> {bindc_name = "a", data_attr = #cuf.cuda<managed>, uniq_name = "_QMm1Fsub1Ea"} -> !fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>
16-
! CHECK: %[[X:.*]] = fir.coordinate_of %[[DT]], x : (!fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
15+
! CHECK: %[[ALLOC:.*]] = cuf.alloc !fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}> {bindc_name = "a", data_attr = #cuf.cuda<managed>, uniq_name = "_QMm1Fsub1Ea"} -> !fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>
16+
! CHECK: %[[DT:.*]]:2 = hlfir.declare %[[ALLOC]] {data_attr = #cuf.cuda<managed>, uniq_name = "_QMm1Fsub1Ea"} : (!fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>) -> (!fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>, !fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>)
17+
! CHECK: fir.address_of(@_QQ_QMm1Tty_device.DerivedInit)
18+
! CHECK: fir.copy
19+
! CHECK: %[[X:.*]] = fir.coordinate_of %[[DT]]#0, x : (!fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
1720
! CHECK: cuf.set_allocator_idx %[[X]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {data_attr = #cuf.cuda<device>}
18-
! CHECK: %[[Z:.*]] = fir.coordinate_of %[[DT]], z : (!fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
21+
! CHECK: %[[Z:.*]] = fir.coordinate_of %[[DT]]#0, z : (!fir.ref<!fir.type<_QMm1Tty_device{x:!fir.box<!fir.heap<!fir.array<?xi32>>>,y:i32,z:!fir.box<!fir.heap<!fir.array<?xi32>>>}>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
1922
! CHECK: cuf.set_allocator_idx %[[Z]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {data_attr = #cuf.cuda<device>}
2023

2124
end module

‎lldb/test/API/commands/frame/var-dil/basics/QualifiedId/TestFrameVarDILQualifiedId.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@ class TestFrameVarDILQualifiedId(TestBase):
1818
# each debug info format.
1919
NO_DEBUG_INFO_TESTCASE = True
2020

21+
@skipIfWindows
2122
def test_frame_var(self):
2223
self.build()
2324
lldbutil.run_to_source_breakpoint(

0 commit comments

Comments
 (0)