Skip to content

Commit f4d87c4

Browse files
clementvalclementval
authored
[flang][cuda] Add asyncId to allocate entry point (#134947)
1 parent fa273e1 commit f4d87c4

38 files changed

+110
-91
lines changed

‎flang-rt/include/flang-rt/runtime/descriptor.h

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,9 @@
2929
#include <cstdio>
3030
#include <cstring>
3131

32+
/// Value used for asyncId when no specific stream is specified.
33+
static constexpr std::int64_t kNoAsyncId = -1;
34+
3235
namespace Fortran::runtime {
3336

3437
class Terminator;
@@ -369,7 +372,7 @@ class Descriptor {
369372
// before calling. It (re)computes the byte strides after
370373
// allocation. Does not allocate automatic components or
371374
// perform default component initialization.
372-
RT_API_ATTRS int Allocate();
375+
RT_API_ATTRS int Allocate(std::int64_t asyncId);
373376
RT_API_ATTRS void SetByteStrides();
374377

375378
// Deallocates storage; does not call FINAL subroutines or

‎flang-rt/include/flang-rt/runtime/reduction-templates.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -347,7 +347,7 @@ inline RT_API_ATTRS void DoMaxMinNorm2(Descriptor &result, const Descriptor &x,
347347
// as the element size of the source.
348348
result.Establish(x.type(), x.ElementBytes(), nullptr, 0, nullptr,
349349
CFI_attribute_allocatable);
350-
if (int stat{result.Allocate()}) {
350+
if (int stat{result.Allocate(kNoAsyncId)}) {
351351
terminator.Crash(
352352
"%s: could not allocate memory for result; STAT=%d", intrinsic, stat);
353353
}

‎flang-rt/lib/cuda/allocatable.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,7 @@ int RTDEF(CUFAllocatableAllocate)(Descriptor &desc, int64_t stream,
5353
}
5454
// Perform the standard allocation.
5555
int stat{RTNAME(AllocatableAllocate)(
56-
desc, hasStat, errMsg, sourceFile, sourceLine)};
56+
desc, stream, hasStat, errMsg, sourceFile, sourceLine)};
5757
if (pinned) {
5858
// Set pinned according to stat. More infrastructre is needed to set it
5959
// closer to the actual allocation call.

‎flang-rt/lib/runtime/allocatable.cpp

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -133,15 +133,17 @@ void RTDEF(AllocatableApplyMold)(
133133
}
134134
}
135135

136-
int RTDEF(AllocatableAllocate)(Descriptor &descriptor, bool hasStat,
137-
const Descriptor *errMsg, const char *sourceFile, int sourceLine) {
136+
int RTDEF(AllocatableAllocate)(Descriptor &descriptor, std::int64_t asyncId,
137+
bool hasStat, const Descriptor *errMsg, const char *sourceFile,
138+
int sourceLine) {
138139
Terminator terminator{sourceFile, sourceLine};
139140
if (!descriptor.IsAllocatable()) {
140141
return ReturnError(terminator, StatInvalidDescriptor, errMsg, hasStat);
141142
} else if (descriptor.IsAllocated()) {
142143
return ReturnError(terminator, StatBaseNotNull, errMsg, hasStat);
143144
} else {
144-
int stat{ReturnError(terminator, descriptor.Allocate(), errMsg, hasStat)};
145+
int stat{
146+
ReturnError(terminator, descriptor.Allocate(asyncId), errMsg, hasStat)};
145147
if (stat == StatOk) {
146148
if (const DescriptorAddendum * addendum{descriptor.Addendum()}) {
147149
if (const auto *derived{addendum->derivedType()}) {
@@ -160,7 +162,7 @@ int RTDEF(AllocatableAllocateSource)(Descriptor &alloc,
160162
const Descriptor &source, bool hasStat, const Descriptor *errMsg,
161163
const char *sourceFile, int sourceLine) {
162164
int stat{RTNAME(AllocatableAllocate)(
163-
alloc, hasStat, errMsg, sourceFile, sourceLine)};
165+
alloc, /*asyncId=*/-1, hasStat, errMsg, sourceFile, sourceLine)};
164166
if (stat == StatOk) {
165167
Terminator terminator{sourceFile, sourceLine};
166168
DoFromSourceAssign(alloc, source, terminator);

‎flang-rt/lib/runtime/array-constructor.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -50,17 +50,17 @@ static RT_API_ATTRS void AllocateOrReallocateVectorIfNeeded(
5050
initialAllocationSize(fromElements, to.ElementBytes())};
5151
to.GetDimension(0).SetBounds(1, allocationSize);
5252
RTNAME(AllocatableAllocate)
53-
(to, /*hasStat=*/false, /*errMsg=*/nullptr, vector.sourceFile,
54-
vector.sourceLine);
53+
(to, /*asyncId=*/-1, /*hasStat=*/false, /*errMsg=*/nullptr,
54+
vector.sourceFile, vector.sourceLine);
5555
to.GetDimension(0).SetBounds(1, fromElements);
5656
vector.actualAllocationSize = allocationSize;
5757
} else {
5858
// Do not over-allocate if the final extent was known before pushing the
5959
// first value: there should be no reallocation.
6060
RUNTIME_CHECK(terminator, previousToElements >= fromElements);
6161
RTNAME(AllocatableAllocate)
62-
(to, /*hasStat=*/false, /*errMsg=*/nullptr, vector.sourceFile,
63-
vector.sourceLine);
62+
(to, /*asyncId=*/-1, /*hasStat=*/false, /*errMsg=*/nullptr,
63+
vector.sourceFile, vector.sourceLine);
6464
vector.actualAllocationSize = previousToElements;
6565
}
6666
} else {

‎flang-rt/lib/runtime/assign.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -99,7 +99,7 @@ static RT_API_ATTRS int AllocateAssignmentLHS(
9999
toDim.SetByteStride(stride);
100100
stride *= toDim.Extent();
101101
}
102-
int result{ReturnError(terminator, to.Allocate())};
102+
int result{ReturnError(terminator, to.Allocate(kNoAsyncId))};
103103
if (result == StatOk && derived && !derived->noInitializationNeeded()) {
104104
result = ReturnError(terminator, Initialize(to, *derived, terminator));
105105
}
@@ -277,7 +277,7 @@ RT_API_ATTRS void Assign(Descriptor &to, const Descriptor &from,
277277
// entity, otherwise, the Deallocate() below will not
278278
// free the descriptor memory.
279279
newFrom.raw().attribute = CFI_attribute_allocatable;
280-
auto stat{ReturnError(terminator, newFrom.Allocate())};
280+
auto stat{ReturnError(terminator, newFrom.Allocate(kNoAsyncId))};
281281
if (stat == StatOk) {
282282
if (HasDynamicComponent(from)) {
283283
// If 'from' has allocatable/automatic component, we cannot

‎flang-rt/lib/runtime/character.cpp

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -117,7 +117,7 @@ static RT_API_ATTRS void Compare(Descriptor &result, const Descriptor &x,
117117
for (int j{0}; j < rank; ++j) {
118118
result.GetDimension(j).SetBounds(1, ub[j]);
119119
}
120-
if (result.Allocate() != CFI_SUCCESS) {
120+
if (result.Allocate(kNoAsyncId) != CFI_SUCCESS) {
121121
terminator.Crash("Compare: could not allocate storage for result");
122122
}
123123
std::size_t xChars{x.ElementBytes() >> shift<CHAR>};
@@ -172,7 +172,7 @@ static RT_API_ATTRS void AdjustLRHelper(Descriptor &result,
172172
for (int j{0}; j < rank; ++j) {
173173
result.GetDimension(j).SetBounds(1, ub[j]);
174174
}
175-
if (result.Allocate() != CFI_SUCCESS) {
175+
if (result.Allocate(kNoAsyncId) != CFI_SUCCESS) {
176176
terminator.Crash("ADJUSTL/R: could not allocate storage for result");
177177
}
178178
for (SubscriptValue resultAt{0}; elements-- > 0;
@@ -226,7 +226,7 @@ static RT_API_ATTRS void LenTrim(Descriptor &result, const Descriptor &string,
226226
for (int j{0}; j < rank; ++j) {
227227
result.GetDimension(j).SetBounds(1, ub[j]);
228228
}
229-
if (result.Allocate() != CFI_SUCCESS) {
229+
if (result.Allocate(kNoAsyncId) != CFI_SUCCESS) {
230230
terminator.Crash("LEN_TRIM: could not allocate storage for result");
231231
}
232232
std::size_t stringElementChars{string.ElementBytes() >> shift<CHAR>};
@@ -408,7 +408,7 @@ static RT_API_ATTRS void GeneralCharFunc(Descriptor &result,
408408
for (int j{0}; j < rank; ++j) {
409409
result.GetDimension(j).SetBounds(1, ub[j]);
410410
}
411-
if (result.Allocate() != CFI_SUCCESS) {
411+
if (result.Allocate(kNoAsyncId) != CFI_SUCCESS) {
412412
terminator.Crash("SCAN/VERIFY: could not allocate storage for result");
413413
}
414414
std::size_t stringElementChars{string.ElementBytes() >> shift<CHAR>};
@@ -511,7 +511,7 @@ static RT_API_ATTRS void MaxMinHelper(Descriptor &accumulator,
511511
for (int j{0}; j < rank; ++j) {
512512
accumulator.GetDimension(j).SetBounds(1, ub[j]);
513513
}
514-
RUNTIME_CHECK(terminator, accumulator.Allocate() == CFI_SUCCESS);
514+
RUNTIME_CHECK(terminator, accumulator.Allocate(kNoAsyncId) == CFI_SUCCESS);
515515
}
516516
for (CHAR *result{accumulator.OffsetElement<CHAR>()}; elements-- > 0;
517517
accumData += accumChars, result += chars, x.IncrementSubscripts(xAt)) {
@@ -587,7 +587,7 @@ void RTDEF(CharacterConcatenate)(Descriptor &accumulator,
587587
for (int j{0}; j < rank; ++j) {
588588
accumulator.GetDimension(j).SetBounds(1, ub[j]);
589589
}
590-
if (accumulator.Allocate() != CFI_SUCCESS) {
590+
if (accumulator.Allocate(kNoAsyncId) != CFI_SUCCESS) {
591591
terminator.Crash(
592592
"CharacterConcatenate: could not allocate storage for result");
593593
}
@@ -610,7 +610,7 @@ void RTDEF(CharacterConcatenateScalar1)(
610610
accumulator.set_base_addr(nullptr);
611611
std::size_t oldLen{accumulator.ElementBytes()};
612612
accumulator.raw().elem_len += chars;
613-
RUNTIME_CHECK(terminator, accumulator.Allocate() == CFI_SUCCESS);
613+
RUNTIME_CHECK(terminator, accumulator.Allocate(kNoAsyncId) == CFI_SUCCESS);
614614
std::memcpy(accumulator.OffsetElement<char>(oldLen), from, chars);
615615
FreeMemory(old);
616616
}
@@ -812,7 +812,7 @@ void RTDEF(Repeat)(Descriptor &result, const Descriptor &string,
812812
std::size_t origBytes{string.ElementBytes()};
813813
result.Establish(string.type(), origBytes * ncopies, nullptr, 0, nullptr,
814814
CFI_attribute_allocatable);
815-
if (result.Allocate() != CFI_SUCCESS) {
815+
if (result.Allocate(kNoAsyncId) != CFI_SUCCESS) {
816816
terminator.Crash("REPEAT could not allocate storage for result");
817817
}
818818
const char *from{string.OffsetElement()};
@@ -846,7 +846,7 @@ void RTDEF(Trim)(Descriptor &result, const Descriptor &string,
846846
}
847847
result.Establish(string.type(), resultBytes, nullptr, 0, nullptr,
848848
CFI_attribute_allocatable);
849-
RUNTIME_CHECK(terminator, result.Allocate() == CFI_SUCCESS);
849+
RUNTIME_CHECK(terminator, result.Allocate(kNoAsyncId) == CFI_SUCCESS);
850850
std::memcpy(result.OffsetElement(), string.OffsetElement(), resultBytes);
851851
}
852852

‎flang-rt/lib/runtime/copy.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -171,7 +171,8 @@ RT_API_ATTRS void CopyElement(const Descriptor &to, const SubscriptValue toAt[],
171171
*reinterpret_cast<Descriptor *>(toPtr + component->offset())};
172172
if (toDesc.raw().base_addr != nullptr) {
173173
toDesc.set_base_addr(nullptr);
174-
RUNTIME_CHECK(terminator, toDesc.Allocate() == CFI_SUCCESS);
174+
RUNTIME_CHECK(
175+
terminator, toDesc.Allocate(/*asyncId=*/-1) == CFI_SUCCESS);
175176
const Descriptor &fromDesc{*reinterpret_cast<const Descriptor *>(
176177
fromPtr + component->offset())};
177178
copyStack.emplace(toDesc, fromDesc);

‎flang-rt/lib/runtime/derived.cpp

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -51,7 +51,8 @@ RT_API_ATTRS int Initialize(const Descriptor &instance,
5151
comp.EstablishDescriptor(allocDesc, instance, terminator);
5252
allocDesc.raw().attribute = CFI_attribute_allocatable;
5353
if (comp.genre() == typeInfo::Component::Genre::Automatic) {
54-
stat = ReturnError(terminator, allocDesc.Allocate(), errMsg, hasStat);
54+
stat = ReturnError(
55+
terminator, allocDesc.Allocate(kNoAsyncId), errMsg, hasStat);
5556
if (stat == StatOk) {
5657
if (const DescriptorAddendum * addendum{allocDesc.Addendum()}) {
5758
if (const auto *derived{addendum->derivedType()}) {
@@ -151,7 +152,8 @@ RT_API_ATTRS int InitializeClone(const Descriptor &clone,
151152
*clone.ElementComponent<Descriptor>(at, comp.offset())};
152153
if (origDesc.IsAllocated()) {
153154
cloneDesc.ApplyMold(origDesc, origDesc.rank());
154-
stat = ReturnError(terminator, cloneDesc.Allocate(), errMsg, hasStat);
155+
stat = ReturnError(
156+
terminator, cloneDesc.Allocate(kNoAsyncId), errMsg, hasStat);
155157
if (stat == StatOk) {
156158
if (const DescriptorAddendum * addendum{cloneDesc.Addendum()}) {
157159
if (const typeInfo::DerivedType *
@@ -258,7 +260,7 @@ static RT_API_ATTRS void CallFinalSubroutine(const Descriptor &descriptor,
258260
copy.raw().attribute = CFI_attribute_allocatable;
259261
Terminator stubTerminator{"CallFinalProcedure() in Fortran runtime", 0};
260262
RUNTIME_CHECK(terminator ? *terminator : stubTerminator,
261-
copy.Allocate() == CFI_SUCCESS);
263+
copy.Allocate(kNoAsyncId) == CFI_SUCCESS);
262264
ShallowCopyDiscontiguousToContiguous(copy, descriptor);
263265
argDescriptor = &copy;
264266
}

‎flang-rt/lib/runtime/descriptor.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -158,7 +158,7 @@ RT_API_ATTRS static inline int MapAllocIdx(const Descriptor &desc) {
158158
#endif
159159
}
160160

161-
RT_API_ATTRS int Descriptor::Allocate() {
161+
RT_API_ATTRS int Descriptor::Allocate(std::int64_t asyncId) {
162162
std::size_t elementBytes{ElementBytes()};
163163
if (static_cast<std::int64_t>(elementBytes) < 0) {
164164
// F'2023 7.4.4.2 p5: "If the character length parameter value evaluates
@@ -170,7 +170,7 @@ RT_API_ATTRS int Descriptor::Allocate() {
170170
// Zero size allocation is possible in Fortran and the resulting
171171
// descriptor must be allocated/associated. Since std::malloc(0)
172172
// result is implementation defined, always allocate at least one byte.
173-
void *p{alloc(byteSize ? byteSize : 1, /*asyncId=*/-1)};
173+
void *p{alloc(byteSize ? byteSize : 1, asyncId)};
174174
if (!p) {
175175
return CFI_ERROR_MEM_ALLOCATION;
176176
}

‎flang-rt/lib/runtime/extrema.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -152,7 +152,7 @@ inline RT_API_ATTRS void CharacterMaxOrMinLoc(const char *intrinsic,
152152
CFI_attribute_allocatable);
153153
result.GetDimension(0).SetBounds(1, extent[0]);
154154
Terminator terminator{source, line};
155-
if (int stat{result.Allocate()}) {
155+
if (int stat{result.Allocate(kNoAsyncId)}) {
156156
terminator.Crash(
157157
"%s: could not allocate memory for result; STAT=%d", intrinsic, stat);
158158
}
@@ -181,7 +181,7 @@ inline RT_API_ATTRS void TotalNumericMaxOrMinLoc(const char *intrinsic,
181181
CFI_attribute_allocatable);
182182
result.GetDimension(0).SetBounds(1, extent[0]);
183183
Terminator terminator{source, line};
184-
if (int stat{result.Allocate()}) {
184+
if (int stat{result.Allocate(kNoAsyncId)}) {
185185
terminator.Crash(
186186
"%s: could not allocate memory for result; STAT=%d", intrinsic, stat);
187187
}

‎flang-rt/lib/runtime/findloc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -220,7 +220,7 @@ void RTDEF(Findloc)(Descriptor &result, const Descriptor &x,
220220
CFI_attribute_allocatable);
221221
result.GetDimension(0).SetBounds(1, extent[0]);
222222
Terminator terminator{source, line};
223-
if (int stat{result.Allocate()}) {
223+
if (int stat{result.Allocate(kNoAsyncId)}) {
224224
terminator.Crash(
225225
"FINDLOC: could not allocate memory for result; STAT=%d", stat);
226226
}

‎flang-rt/lib/runtime/matmul-transpose.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -183,7 +183,7 @@ inline static RT_API_ATTRS void DoMatmulTranspose(
183183
for (int j{0}; j < resRank; ++j) {
184184
result.GetDimension(j).SetBounds(1, extent[j]);
185185
}
186-
if (int stat{result.Allocate()}) {
186+
if (int stat{result.Allocate(kNoAsyncId)}) {
187187
terminator.Crash(
188188
"MATMUL-TRANSPOSE: could not allocate memory for result; STAT=%d",
189189
stat);

‎flang-rt/lib/runtime/matmul.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -255,7 +255,7 @@ static inline RT_API_ATTRS void DoMatmul(
255255
for (int j{0}; j < resRank; ++j) {
256256
result.GetDimension(j).SetBounds(1, extent[j]);
257257
}
258-
if (int stat{result.Allocate()}) {
258+
if (int stat{result.Allocate(kNoAsyncId)}) {
259259
terminator.Crash(
260260
"MATMUL: could not allocate memory for result; STAT=%d", stat);
261261
}

‎flang-rt/lib/runtime/misc-intrinsic.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ static RT_API_ATTRS void TransferImpl(Descriptor &result,
3030
if (const DescriptorAddendum * addendum{mold.Addendum()}) {
3131
*result.Addendum() = *addendum;
3232
}
33-
if (int stat{result.Allocate()}) {
33+
if (int stat{result.Allocate(kNoAsyncId)}) {
3434
Terminator{sourceFile, line}.Crash(
3535
"TRANSFER: could not allocate memory for result; STAT=%d", stat);
3636
}

‎flang-rt/lib/runtime/temporary-stack.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,7 @@ void DescriptorStorage<COPY_VALUES>::push(const Descriptor &source) {
148148
if constexpr (COPY_VALUES) {
149149
// copy the data pointed to by the box
150150
box.set_base_addr(nullptr);
151-
box.Allocate();
151+
box.Allocate(kNoAsyncId);
152152
RTNAME(AssignTemporary)
153153
(box, source, terminator_.sourceFileName(), terminator_.sourceLine());
154154
}

‎flang-rt/lib/runtime/tools.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -261,7 +261,7 @@ RT_API_ATTRS void CreatePartialReductionResult(Descriptor &result,
261261
for (int j{0}; j + 1 < xRank; ++j) {
262262
result.GetDimension(j).SetBounds(1, resultExtent[j]);
263263
}
264-
if (int stat{result.Allocate()}) {
264+
if (int stat{result.Allocate(kNoAsyncId)}) {
265265
terminator.Crash(
266266
"%s: could not allocate memory for result; STAT=%d", intrinsic, stat);
267267
}

‎flang-rt/lib/runtime/transformational.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -132,7 +132,7 @@ static inline RT_API_ATTRS std::size_t AllocateResult(Descriptor &result,
132132
for (int j{0}; j < rank; ++j) {
133133
result.GetDimension(j).SetBounds(1, extent[j]);
134134
}
135-
if (int stat{result.Allocate()}) {
135+
if (int stat{result.Allocate(kNoAsyncId)}) {
136136
terminator.Crash(
137137
"%s: Could not allocate memory for result (stat=%d)", function, stat);
138138
}
@@ -157,7 +157,7 @@ static inline RT_API_ATTRS std::size_t AllocateBesselResult(Descriptor &result,
157157
for (int j{0}; j < rank; ++j) {
158158
result.GetDimension(j).SetBounds(1, extent[j]);
159159
}
160-
if (int stat{result.Allocate()}) {
160+
if (int stat{result.Allocate(kNoAsyncId)}) {
161161
terminator.Crash(
162162
"%s: Could not allocate memory for result (stat=%d)", function, stat);
163163
}

‎flang-rt/unittests/Evaluate/reshape.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ int main() {
2626
for (int j{0}; j < 3; ++j) {
2727
source->GetDimension(j).SetBounds(1, sourceExtent[j]);
2828
}
29-
TEST(source->Allocate() == CFI_SUCCESS);
29+
TEST(source->Allocate(kNoAsyncId) == CFI_SUCCESS);
3030
TEST(source->IsAllocated());
3131
MATCH(2, source->GetDimension(0).Extent());
3232
MATCH(3, source->GetDimension(1).Extent());

‎flang-rt/unittests/Runtime/Allocatable.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ TEST(AllocatableTest, MoveAlloc) {
2626
auto b{createAllocatable(TypeCategory::Integer, 4)};
2727
// ALLOCATE(a(20))
2828
a->GetDimension(0).SetBounds(1, 20);
29-
a->Allocate();
29+
a->Allocate(kNoAsyncId);
3030

3131
EXPECT_TRUE(a->IsAllocated());
3232
EXPECT_FALSE(b->IsAllocated());
@@ -46,7 +46,7 @@ TEST(AllocatableTest, MoveAlloc) {
4646
// move_alloc with errMsg
4747
auto errMsg{Descriptor::Create(
4848
sizeof(char), 64, nullptr, 0, nullptr, CFI_attribute_allocatable)};
49-
errMsg->Allocate();
49+
errMsg->Allocate(kNoAsyncId);
5050
RTNAME(MoveAlloc)(*b, *a, nullptr, false, errMsg.get(), __FILE__, __LINE__);
5151
EXPECT_FALSE(a->IsAllocated());
5252
EXPECT_TRUE(b->IsAllocated());

‎flang-rt/unittests/Runtime/CUDA/Allocatable.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ TEST(AllocatableCUFTest, SimpleDeviceAllocatable) {
4242
CUDA_REPORT_IF_ERROR(cudaMalloc(&device_desc, a->SizeInBytes()));
4343

4444
RTNAME(AllocatableAllocate)
45-
(*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
45+
(*a, kNoAsyncId, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
4646
EXPECT_TRUE(a->IsAllocated());
4747
RTNAME(CUFDescriptorSync)(device_desc, a.get(), __FILE__, __LINE__);
4848
cudaDeviceSynchronize();

‎flang-rt/unittests/Runtime/CUDA/AllocatorCUF.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,8 @@ TEST(AllocatableCUFTest, SimpleDeviceAllocate) {
3535
EXPECT_FALSE(a->HasAddendum());
3636
RTNAME(AllocatableSetBounds)(*a, 0, 1, 10);
3737
RTNAME(AllocatableAllocate)
38-
(*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
38+
(*a, /*asyncId=*/-1, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__,
39+
__LINE__);
3940
EXPECT_TRUE(a->IsAllocated());
4041
RTNAME(AllocatableDeallocate)
4142
(*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
@@ -53,7 +54,8 @@ TEST(AllocatableCUFTest, SimplePinnedAllocate) {
5354
EXPECT_FALSE(a->HasAddendum());
5455
RTNAME(AllocatableSetBounds)(*a, 0, 1, 10);
5556
RTNAME(AllocatableAllocate)
56-
(*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
57+
(*a, /*asyncId=*/-1, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__,
58+
__LINE__);
5759
EXPECT_TRUE(a->IsAllocated());
5860
RTNAME(AllocatableDeallocate)
5961
(*a, /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);

0 commit comments

Comments
 (0)