[flang-commits] [flang] [llvm] [Flang-RT] Change alignment for allocatable and pointer arrays to 64 … (PR #206525)
Jason Van Beusekom via flang-commits
flang-commits at lists.llvm.org
Mon Jun 29 09:53:48 PDT 2026
https://github.com/Jason-Van-Beusekom created https://github.com/llvm/llvm-project/pull/206525
…bytes
in https://github.com/llvm/llvm-project/pull/194969 the default array alignment was changed to 64bytes for Fortran globals. In this PR the malloc wrapper in the flang runtime is modified to accept an alignment argument (0 default) and for cases above 16 byte (std::max_align_t) uses `aligned_alloc` instead of malloc.
The issue of alignment was discussed in this RFC: https://discourse.llvm.org/t/rfc-alignment-of-global-arrays/90397/13
Assisted with Opus 4.6
>From e174b89bf38e9949ad2dec2d544b8e5f510a2e97 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Mon, 29 Jun 2026 09:46:03 -0500
Subject: [PATCH] [Flang-RT] Change alignment for allocatable and pointer
arrays to 64 bytes
---
.../flang-rt/runtime/allocator-registry.h | 17 +++++++++--
flang-rt/lib/cuda/allocator.cpp | 20 ++++++++-----
flang-rt/lib/cuda/descriptor.cpp | 2 +-
flang-rt/lib/runtime/ISO_Fortran_binding.cpp | 4 ++-
flang-rt/lib/runtime/descriptor.cpp | 4 ++-
flang-rt/lib/runtime/pointer.cpp | 8 +++--
flang-rt/unittests/Runtime/Allocatable.cpp | 25 ++++++++++++++++
flang-rt/unittests/Runtime/Pointer.cpp | 29 +++++++++++++++++++
flang/include/flang/Runtime/CUDA/allocator.h | 8 ++---
.../flang/Runtime/allocator-registry-consts.h | 1 +
flang/include/flang/Runtime/pointer.h | 2 +-
11 files changed, 98 insertions(+), 22 deletions(-)
diff --git a/flang-rt/include/flang-rt/runtime/allocator-registry.h b/flang-rt/include/flang-rt/runtime/allocator-registry.h
index f0ba77a360736..d2ba6b109b161 100644
--- a/flang-rt/include/flang-rt/runtime/allocator-registry.h
+++ b/flang-rt/include/flang-rt/runtime/allocator-registry.h
@@ -19,7 +19,7 @@
namespace Fortran::runtime {
-using AllocFct = void *(*)(std::size_t, std::int64_t *);
+using AllocFct = void *(*)(std::size_t, std::size_t, std::int64_t *);
using FreeFct = void (*)(void *);
typedef struct Allocator_t {
@@ -27,8 +27,19 @@ typedef struct Allocator_t {
FreeFct free{nullptr};
} Allocator_t;
-static RT_API_ATTRS void *MallocWrapper(
- std::size_t size, [[maybe_unused]] std::int64_t *) {
+static RT_API_ATTRS void *MallocWrapper(std::size_t size,
+ [[maybe_unused]] std::size_t alignment, [[maybe_unused]] std::int64_t *) {
+#if !defined(RT_DEVICE_COMPILATION) && !defined(_WIN32)
+ // std::malloc only guarantees alignof(std::max_align_t). When a larger
+ // alignment is requested, use std::aligned_alloc.
+ if (alignment > alignof(std::max_align_t)) {
+ // Round size up to a multiple of the alignment
+ if (size > SIZE_MAX - alignment)
+ return nullptr;
+ std::size_t alignedSize{((size + alignment - 1) / alignment) * alignment};
+ return std::aligned_alloc(alignment, alignedSize);
+ }
+#endif
return std::malloc(size);
}
#ifdef RT_DEVICE_COMPILATION
diff --git a/flang-rt/lib/cuda/allocator.cpp b/flang-rt/lib/cuda/allocator.cpp
index 8bc32678fe568..a0214c6004de2 100644
--- a/flang-rt/lib/cuda/allocator.cpp
+++ b/flang-rt/lib/cuda/allocator.cpp
@@ -164,8 +164,9 @@ int RTDECL(CUFSetAssociatedStream)(void *p, cudaStream_t stream) {
}
}
-void *CUFAllocPinned(
- std::size_t sizeInBytes, [[maybe_unused]] std::int64_t *asyncObject) {
+void *CUFAllocPinned(std::size_t sizeInBytes,
+ [[maybe_unused]] std::size_t alignment,
+ [[maybe_unused]] std::int64_t *asyncObject) {
void *p;
CUDA_REPORT_IF_ERROR(cudaMallocHost((void **)&p, sizeInBytes));
return p;
@@ -173,7 +174,8 @@ void *CUFAllocPinned(
void CUFFreePinned(void *p) { cudaFreeHost(p); }
-void *CUFAllocDevice(std::size_t sizeInBytes, std::int64_t *asyncObject) {
+void *CUFAllocDevice(std::size_t sizeInBytes,
+ [[maybe_unused]] std::size_t alignment, std::int64_t *asyncObject) {
void *p;
if (Fortran::runtime::executionEnvironment.cudaDeviceIsManaged) {
CUDA_REPORT_IF_ERROR(
@@ -202,8 +204,9 @@ void CUFFreeDevice(void *p) {
}
}
-void *CUFAllocManaged(
- std::size_t sizeInBytes, [[maybe_unused]] std::int64_t *asyncObject) {
+void *CUFAllocManaged(std::size_t sizeInBytes,
+ [[maybe_unused]] std::size_t alignment,
+ [[maybe_unused]] std::int64_t *asyncObject) {
void *p;
CUDA_REPORT_IF_ERROR(
cudaMallocManaged((void **)&p, sizeInBytes, cudaMemAttachGlobal));
@@ -212,10 +215,11 @@ void *CUFAllocManaged(
void CUFFreeManaged(void *p) { CUDA_REPORT_IF_ERROR(cudaFree(p)); }
-void *CUFAllocUnified(
- std::size_t sizeInBytes, [[maybe_unused]] std::int64_t *asyncObject) {
+void *CUFAllocUnified(std::size_t sizeInBytes,
+ [[maybe_unused]] std::size_t alignment,
+ [[maybe_unused]] std::int64_t *asyncObject) {
// Call alloc managed for the time being.
- return CUFAllocManaged(sizeInBytes, asyncObject);
+ return CUFAllocManaged(sizeInBytes, alignment, asyncObject);
}
void CUFFreeUnified(void *p) {
diff --git a/flang-rt/lib/cuda/descriptor.cpp b/flang-rt/lib/cuda/descriptor.cpp
index 2cf795181ea7b..1af62ca3619e6 100644
--- a/flang-rt/lib/cuda/descriptor.cpp
+++ b/flang-rt/lib/cuda/descriptor.cpp
@@ -21,7 +21,7 @@ RT_EXT_API_GROUP_BEGIN
Descriptor *RTDEF(CUFAllocDescriptor)(
std::size_t sizeInBytes, const char *sourceFile, int sourceLine) {
return reinterpret_cast<Descriptor *>(
- CUFAllocManaged(sizeInBytes, /*asyncObject=*/nullptr));
+ CUFAllocManaged(sizeInBytes, /*alignment=*/0, /*asyncObject=*/nullptr));
}
void RTDEF(CUFFreeDescriptor)(
diff --git a/flang-rt/lib/runtime/ISO_Fortran_binding.cpp b/flang-rt/lib/runtime/ISO_Fortran_binding.cpp
index a5f8b357ae0b8..39f409cf62788 100644
--- a/flang-rt/lib/runtime/ISO_Fortran_binding.cpp
+++ b/flang-rt/lib/runtime/ISO_Fortran_binding.cpp
@@ -76,7 +76,9 @@ RT_API_ATTRS int CFI_allocate(CFI_cdesc_t *descriptor,
dim->sm = byteSize;
byteSize *= extent;
}
- void *p{runtime::AllocateValidatedPointerPayload(byteSize)};
+ std::size_t alignment{rank > 0 ? kDefaultArrayAlignment : 0};
+ void *p{runtime::AllocateValidatedPointerPayload(
+ byteSize, /*allocatorIdx=*/0, alignment)};
if (!p && byteSize) {
return CFI_ERROR_MEM_ALLOCATION;
}
diff --git a/flang-rt/lib/runtime/descriptor.cpp b/flang-rt/lib/runtime/descriptor.cpp
index 6c9e76afb117e..4aacbf6ae7949 100644
--- a/flang-rt/lib/runtime/descriptor.cpp
+++ b/flang-rt/lib/runtime/descriptor.cpp
@@ -174,6 +174,8 @@ RT_API_ATTRS int Descriptor::Allocate(std::int64_t *asyncObject) {
}
std::size_t byteSize{Elements() * elementBytes};
AllocFct alloc{allocatorRegistry.GetAllocator(MapAllocIdx())};
+ // For arrays, use the default array alignment
+ std::size_t alignment{rank() > 0 ? kDefaultArrayAlignment : 0};
// Zero size allocation is possible in Fortran and the resulting
// descriptor must be allocated/associated. Since std::malloc(0)
// result is implementation defined, always allocate at least one byte.
@@ -183,7 +185,7 @@ RT_API_ATTRS int Descriptor::Allocate(std::int64_t *asyncObject) {
}
byteSize = 1;
}
- void *p{alloc(byteSize, asyncObject)};
+ void *p{alloc(byteSize, alignment, asyncObject)};
if (!p) {
return CFI_ERROR_MEM_ALLOCATION;
}
diff --git a/flang-rt/lib/runtime/pointer.cpp b/flang-rt/lib/runtime/pointer.cpp
index 0832b5656f1ab..a94ba4aa838cf 100644
--- a/flang-rt/lib/runtime/pointer.cpp
+++ b/flang-rt/lib/runtime/pointer.cpp
@@ -138,13 +138,13 @@ void RTDEF(PointerAssociateRemappingMonomorphic)(Descriptor &pointer,
}
RT_API_ATTRS void *AllocateValidatedPointerPayload(
- std::size_t byteSize, int allocatorIdx) {
+ std::size_t byteSize, int allocatorIdx, std::size_t alignment) {
// Add space for a footer to validate during deallocation.
constexpr std::size_t align{sizeof(std::uintptr_t)};
byteSize = ((byteSize + align - 1) / align) * align;
std::size_t total{byteSize + sizeof(std::uintptr_t)};
AllocFct alloc{allocatorRegistry.GetAllocator(allocatorIdx)};
- void *p{alloc(total, /*asyncObject=*/nullptr)};
+ void *p{alloc(total, alignment, /*asyncObject=*/nullptr)};
if (p && allocatorIdx == 0) {
// Fill the footer word with the XOR of the ones' complement of
// the base address, which is a value that would be highly unlikely
@@ -170,7 +170,9 @@ int RTDEF(PointerAllocate)(Descriptor &pointer, bool hasStat,
elementBytes = pointer.raw().elem_len = 0;
}
std::size_t byteSize{pointer.Elements() * elementBytes};
- void *p{AllocateValidatedPointerPayload(byteSize, pointer.GetAllocIdx())};
+ std::size_t alignment{pointer.rank() > 0 ? kDefaultArrayAlignment : 0};
+ void *p{AllocateValidatedPointerPayload(
+ byteSize, pointer.GetAllocIdx(), alignment)};
if (!p) {
return ReturnError(terminator, CFI_ERROR_MEM_ALLOCATION, errMsg, hasStat);
}
diff --git a/flang-rt/unittests/Runtime/Allocatable.cpp b/flang-rt/unittests/Runtime/Allocatable.cpp
index b394312e5bc5a..28852362e10a6 100644
--- a/flang-rt/unittests/Runtime/Allocatable.cpp
+++ b/flang-rt/unittests/Runtime/Allocatable.cpp
@@ -75,6 +75,31 @@ TEST(AllocatableTest, MoveAlloc) {
EXPECT_EQ(errStr, "MOVE_ALLOC passed the same address as to and from");
}
+TEST(AllocatableTest, AllocateArrayIsAligned) {
+ using Fortran::common::TypeCategory;
+#if !defined(_WIN32)
+ // Since a single std::malloc result may happen to be 64-byte aligned
+ // allocate multiple arrays and check that all of them are aligned correctly
+ constexpr int count{32};
+ OwningPtr<Descriptor> arrays[count];
+ for (int i{0}; i < count; ++i) {
+ // REAL(8), ALLOCATABLE :: a(:)
+ arrays[i] = createAllocatable(TypeCategory::Real, 8);
+ // ALLOCATE(a(100))
+ arrays[i]->GetDimension(0).SetBounds(1, 100);
+ arrays[i]->Allocate(kNoAsyncObject);
+ EXPECT_TRUE(arrays[i]->IsAllocated());
+ EXPECT_EQ(reinterpret_cast<std::uintptr_t>(arrays[i]->raw().base_addr) %
+ kDefaultArrayAlignment,
+ 0u);
+ }
+ for (int i{0}; i < count; ++i) {
+ arrays[i]->Deallocate();
+ EXPECT_FALSE(arrays[i]->IsAllocated());
+ }
+#endif
+}
+
TEST(AllocatableTest, AllocateFromScalarSource) {
using Fortran::common::TypeCategory;
// REAL(4), ALLOCATABLE :: a(:)
diff --git a/flang-rt/unittests/Runtime/Pointer.cpp b/flang-rt/unittests/Runtime/Pointer.cpp
index 6e8861d374889..da1c07f59df6f 100644
--- a/flang-rt/unittests/Runtime/Pointer.cpp
+++ b/flang-rt/unittests/Runtime/Pointer.cpp
@@ -31,6 +31,35 @@ TEST(Pointer, BasicAllocateDeallocate) {
EXPECT_FALSE(RTNAME(PointerIsAssociated)(*p));
}
+TEST(Pointer, AllocateArrayIsAligned) {
+#if !defined(_WIN32)
+ // Since a single std::malloc result may happen to be 64-byte aligned
+ // allocate multiple arrays and check that all of them are aligned correctly
+ constexpr int count{32};
+ OwningPtr<Descriptor> pointers[count];
+ for (int i{0}; i < count; ++i) {
+ // REAL(8), POINTER :: p(:)
+ pointers[i] = Descriptor::Create(
+ TypeCode{Fortran::common::TypeCategory::Real, 8}, 8, nullptr, 1, nullptr,
+ CFI_attribute_pointer);
+ // ALLOCATE(p(100))
+ RTNAME(PointerSetBounds)(*pointers[i], 0, 1, 100);
+ RTNAME(PointerAllocate)
+ (*pointers[i], /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
+ EXPECT_TRUE(RTNAME(PointerIsAssociated)(*pointers[i]));
+ EXPECT_EQ(reinterpret_cast<std::uintptr_t>(pointers[i]->raw().base_addr) %
+ kDefaultArrayAlignment,
+ 0u);
+ }
+ for (int i{0}; i < count; ++i) {
+ // DEALLOCATE(p)
+ RTNAME(PointerDeallocate)
+ (*pointers[i], /*hasStat=*/false, /*errMsg=*/nullptr, __FILE__, __LINE__);
+ EXPECT_FALSE(RTNAME(PointerIsAssociated)(*pointers[i]));
+ }
+#endif
+}
+
TEST(Pointer, ApplyMoldAllocation) {
// REAL(4), POINTER :: p
auto m{Descriptor::Create(TypeCode{Fortran::common::TypeCategory::Real, 4}, 4,
diff --git a/flang/include/flang/Runtime/CUDA/allocator.h b/flang/include/flang/Runtime/CUDA/allocator.h
index 18907252b575d..8c1e0f0f4f690 100644
--- a/flang/include/flang/Runtime/CUDA/allocator.h
+++ b/flang/include/flang/Runtime/CUDA/allocator.h
@@ -25,16 +25,16 @@ void RTDECL(CUFRegisterAllocator)();
void CUFResetStream(cudaStream_t stream);
-void *CUFAllocPinned(std::size_t, std::int64_t *);
+void *CUFAllocPinned(std::size_t, std::size_t, std::int64_t *);
void CUFFreePinned(void *);
-void *CUFAllocDevice(std::size_t, std::int64_t *);
+void *CUFAllocDevice(std::size_t, std::size_t, std::int64_t *);
void CUFFreeDevice(void *);
-void *CUFAllocManaged(std::size_t, std::int64_t *);
+void *CUFAllocManaged(std::size_t, std::size_t, std::int64_t *);
void CUFFreeManaged(void *);
-void *CUFAllocUnified(std::size_t, std::int64_t *);
+void *CUFAllocUnified(std::size_t, std::size_t, std::int64_t *);
void CUFFreeUnified(void *);
} // namespace Fortran::runtime::cuda
diff --git a/flang/include/flang/Runtime/allocator-registry-consts.h b/flang/include/flang/Runtime/allocator-registry-consts.h
index a5f52749e945d..d8e8274d77ae9 100644
--- a/flang/include/flang/Runtime/allocator-registry-consts.h
+++ b/flang/include/flang/Runtime/allocator-registry-consts.h
@@ -18,6 +18,7 @@ static constexpr unsigned kPinnedAllocatorPos = 1;
static constexpr unsigned kDeviceAllocatorPos = 2;
static constexpr unsigned kManagedAllocatorPos = 3;
static constexpr unsigned kUnifiedAllocatorPos = 4;
+static constexpr unsigned kDefaultArrayAlignment = 64;
RT_OFFLOAD_VAR_GROUP_END
diff --git a/flang/include/flang/Runtime/pointer.h b/flang/include/flang/Runtime/pointer.h
index d05bc5276c295..832cb169c19b6 100644
--- a/flang/include/flang/Runtime/pointer.h
+++ b/flang/include/flang/Runtime/pointer.h
@@ -130,7 +130,7 @@ bool RTDECL(PointerIsAssociatedWith)(
// Fortran POINTERs are allocated with an extra validation word after their
// payloads in order to detect erroneous deallocations later.
RT_API_ATTRS void *AllocateValidatedPointerPayload(
- std::size_t, int allocatorIdx = 0);
+ std::size_t, int allocatorIdx = 0, std::size_t alignment = 0);
RT_API_ATTRS bool ValidatePointerPayload(const ISO::CFI_cdesc_t &);
} // extern "C"
More information about the flang-commits
mailing list