Skip to content

Conversation

skc7
Copy link
Contributor

@skc7 skc7 commented Aug 8, 2025

This PR adds below APIs to flang-rt:
DescriptorGetBaseAddress to retrive base_addr from Descriptor
DescriptorGetDataSizeInBytes to retrive the total Size in bytes of data from Descriptor.

@skc7 skc7 requested a review from mjklemm August 8, 2025 16:35
@mjklemm mjklemm requested review from kparzysz and ergawy August 14, 2025 12:04
Copy link
Contributor

@mjklemm mjklemm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@skc7 skc7 marked this pull request as ready for review August 18, 2025 13:35
@llvmbot llvmbot added flang Flang issues not falling into any other category flang:fir-hlfir labels Aug 18, 2025
@llvmbot
Copy link
Member

llvmbot commented Aug 18, 2025

@llvm/pr-subscribers-flang-fir-hlfir

Author: Chaitanya (skc7)

Changes

This PR adds below APIs to flang-rt:
DescriptorGetBaseAddress to retrive base_addr from Descriptor
DescriptorGetDataSizeInBytes to retrive the total Size in bytes of data from Descriptor.


Full diff: https://github.com/llvm/llvm-project/pull/152756.diff

5 Files Affected:

  • (modified) flang-rt/lib/runtime/support.cpp (+24)
  • (modified) flang-rt/unittests/Runtime/Support.cpp (+23)
  • (modified) flang/include/flang/Optimizer/Builder/Runtime/Support.h (+13)
  • (modified) flang/include/flang/Runtime/support.h (+8)
  • (modified) flang/lib/Optimizer/Builder/Runtime/Support.cpp (+21)
diff --git a/flang-rt/lib/runtime/support.cpp b/flang-rt/lib/runtime/support.cpp
index 9beb46e48a11e..ffeaafaa162ea 100644
--- a/flang-rt/lib/runtime/support.cpp
+++ b/flang-rt/lib/runtime/support.cpp
@@ -48,6 +48,30 @@ void RTDEF(CopyAndUpdateDescriptor)(Descriptor &to, const Descriptor &from,
   }
 }
 
+void *RTDEF(DescriptorGetBaseAddress)(
+    const Descriptor &desc, const char *sourceFile, int sourceLine) {
+  Terminator terminator{sourceFile, sourceLine};
+  void *baseAddr = desc.raw().base_addr;
+  if (!baseAddr) {
+    terminator.Crash("Could not retrieve Descriptor's base address");
+  }
+  return baseAddr;
+}
+
+std::size_t RTDEF(DescriptorGetDataSizeInBytes)(
+    const Descriptor &desc, const char *sourceFile, int sourceLine) {
+  Terminator terminator{sourceFile, sourceLine};
+  std::size_t descElements{desc.Elements()};
+  if (!descElements) {
+    terminator.Crash("Could not retrieve Descriptor's Elements");
+  }
+  std::size_t descElementBytes{desc.ElementBytes()};
+  if (!descElementBytes) {
+    terminator.Crash("Could not retrieve Descriptor's ElementBytes");
+  }
+  return descElements * descElementBytes;
+}
+
 RT_EXT_API_GROUP_END
 } // extern "C"
 } // namespace Fortran::runtime
diff --git a/flang-rt/unittests/Runtime/Support.cpp b/flang-rt/unittests/Runtime/Support.cpp
index 46c6805d5d238..264dde872c242 100644
--- a/flang-rt/unittests/Runtime/Support.cpp
+++ b/flang-rt/unittests/Runtime/Support.cpp
@@ -98,3 +98,26 @@ TEST(IsContiguous, Basic) {
   EXPECT_TRUE(RTNAME(IsContiguousUpTo)(section, 1));
   EXPECT_FALSE(RTNAME(IsContiguousUpTo)(section, 2));
 }
+
+TEST(DescriptorGetBaseAddress, Basic) {
+  auto array{MakeArray<TypeCategory::Integer, 4>(
+      std::vector<int>{2, 3}, std::vector<std::int32_t>{0, 1, 2, 3, 4, 5})};
+  void *baseAddr = RTNAME(DescriptorGetBaseAddress)(*array);
+  EXPECT_NE(baseAddr, nullptr);
+  EXPECT_EQ(baseAddr, array->raw().base_addr);
+}
+
+TEST(DescriptorGetDataSizeInBytes, Basic) {
+  // Test with a 2x3 integer*4 array
+  auto int4Array{MakeArray<TypeCategory::Integer, 4>({2, 3})};
+  EXPECT_EQ(RTNAME(DescriptorGetDataSizeInBytes)(*int4Array),
+      6 * sizeof(std::int32_t));
+  // Test with a 1D, 5-element real*8 array
+  auto real8Array{MakeArray<TypeCategory::Real, 8>({5})};
+  EXPECT_EQ(
+      RTNAME(DescriptorGetDataSizeInBytes)(*real8Array), 5 * sizeof(double));
+  // Test with a scalar logical*1
+  auto logical1Scalar{MakeArray<TypeCategory::Logical, 1>({})};
+  EXPECT_EQ(
+      RTNAME(DescriptorGetDataSizeInBytes)(*logical1Scalar), 1 * sizeof(bool));
+}
diff --git a/flang/include/flang/Optimizer/Builder/Runtime/Support.h b/flang/include/flang/Optimizer/Builder/Runtime/Support.h
index d0a474d75d2eb..41db61c19b07e 100644
--- a/flang/include/flang/Optimizer/Builder/Runtime/Support.h
+++ b/flang/include/flang/Optimizer/Builder/Runtime/Support.h
@@ -31,5 +31,18 @@ void genCopyAndUpdateDescriptor(fir::FirOpBuilder &builder, mlir::Location loc,
 mlir::Value genIsAssumedSize(fir::FirOpBuilder &builder, mlir::Location loc,
                              mlir::Value box);
 
+/// Generate call to `DescriptorGetBaseAddress` runtime routine.
+mlir::Value genDescriptorGetBaseAddress(fir::FirOpBuilder &builder,
+                                        mlir::Location loc, mlir::Value desc,
+                                        mlir::Value sourceFile,
+                                        mlir::Value sourceLine);
+
+/// Generate call to `DescriptorGetDataSizeInBytes` runtime routine.
+mlir::Value genDescriptorGetDataSizeInBytes(fir::FirOpBuilder &builder,
+                                            mlir::Location loc,
+                                            mlir::Value desc,
+                                            mlir::Value sourceFile,
+                                            mlir::Value sourceLine);
+
 } // namespace fir::runtime
 #endif // FORTRAN_OPTIMIZER_BUILDER_RUNTIME_SUPPORT_H
diff --git a/flang/include/flang/Runtime/support.h b/flang/include/flang/Runtime/support.h
index 8a345bee7f867..5ebe6c6406a01 100644
--- a/flang/include/flang/Runtime/support.h
+++ b/flang/include/flang/Runtime/support.h
@@ -49,6 +49,14 @@ void RTDECL(CopyAndUpdateDescriptor)(Descriptor &to, const Descriptor &from,
     const typeInfo::DerivedType *newDynamicType,
     ISO::CFI_attribute_t newAttribute, enum LowerBoundModifier newLowerBounds);
 
+// Retrieve the base_addr from Descriptor
+void *RTDECL(DescriptorGetBaseAddress)(const Descriptor &desc,
+    const char *sourceFile = nullptr, int sourceLine = 0);
+
+// Retrieve the totalSizeInBytes of data from Descriptor
+std::size_t RTDECL(DescriptorGetDataSizeInBytes)(const Descriptor &desc,
+    const char *sourceFile = nullptr, int sourceLine = 0);
+
 } // extern "C"
 } // namespace Fortran::runtime
 #endif // FORTRAN_RUNTIME_SUPPORT_H_
diff --git a/flang/lib/Optimizer/Builder/Runtime/Support.cpp b/flang/lib/Optimizer/Builder/Runtime/Support.cpp
index d0d48ad718da4..12994b596df4b 100644
--- a/flang/lib/Optimizer/Builder/Runtime/Support.cpp
+++ b/flang/lib/Optimizer/Builder/Runtime/Support.cpp
@@ -54,3 +54,24 @@ mlir::Value fir::runtime::genIsAssumedSize(fir::FirOpBuilder &builder,
   auto args = fir::runtime::createArguments(builder, loc, fTy, box);
   return fir::CallOp::create(builder, loc, func, args).getResult(0);
 }
+
+mlir::Value fir::runtime::genDescriptorGetBaseAddress(
+    fir::FirOpBuilder &builder, mlir::Location loc, mlir::Value desc,
+    mlir::Value sourceFile, mlir::Value sourceLine) {
+  mlir::func::FuncOp baseAddrFunc =
+      fir::runtime::getRuntimeFunc<mkRTKey(DescriptorGetBaseAddress)>(loc,
+                                                                      builder);
+  llvm::SmallVector<mlir::Value> args{desc, sourceFile, sourceLine};
+  return fir::CallOp::create(builder, loc, baseAddrFunc, args).getResult(0);
+}
+
+mlir::Value fir::runtime::genDescriptorGetDataSizeInBytes(
+    fir::FirOpBuilder &builder, mlir::Location loc, mlir::Value desc,
+    mlir::Value sourceFile, mlir::Value sourceLine) {
+  mlir::func::FuncOp getDataSizeInBytesFunc =
+      fir::runtime::getRuntimeFunc<mkRTKey(DescriptorGetDataSizeInBytes)>(
+          loc, builder);
+  llvm::SmallVector<mlir::Value> args{desc, sourceFile, sourceLine};
+  return fir::CallOp::create(builder, loc, getDataSizeInBytesFunc, args)
+      .getResult(0);
+}

Copy link
Contributor

@clementval clementval left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the use case for that? We have an operation for the base addr already.

Copy link
Contributor

@vzakhari vzakhari left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please answer Valentin's questions before we proceed with this PR.

@skc7
Copy link
Contributor Author

skc7 commented Aug 19, 2025

@llvm/pr-subscribers-flang-fir-hlfir

Author: Chaitanya (skc7)

Changes

Hi @clementval

This PR is pre-requisite for #140523
Have to do omp_target_memcpy between two device ptrs, and to do that, need to access the base_addr and DataSizeInBytes from descriptor.
I'm not aware of exisiting API regarding base_addr, If its already present, I will modify this PR accordingly.

@clementval
Copy link
Contributor

There is a fir.box_addr operation that will return the base address directly without the need of a runtime call. But in your case I think you probably want to have a specific entry point for OpenMP that takes two flang descriptors or a flang descriptor and a pointer because you will need to take into account more than just the base adds and the size. You need to know if both descriptor are contiguous or not, if they are the same size and so on.

There is a pretty similar use case with the CUDA Fortran data transfer on assignment and the runtime entry point is taking descriptor directly. I would suggest to do the same.

@skc7
Copy link
Contributor Author

skc7 commented Aug 19, 2025

There is a fir.box_addr operation that will return the base address directly without the need of a runtime call. But in your case I think you probably want to have a specific entry point for OpenMP that takes two flang descriptors or a flang descriptor and a pointer because you will need to take into account more than just the base adds and the size. You need to know if both descriptor are contiguous or not, if they are the same size and so on.

There is a pretty similar use case with the CUDA Fortran data transfer on assignment and the runtime entry point is taking descriptor directly. I would suggest to do the same.

For PR #140523, fortran runtime assign call in omp.target needs to be hoisted out, and to be replaced by new API, which does memcpy of device pointers. But this required openmp runtime dependency as in #145465(initial PR for this work).

I have looked at CUDA flang-rt APIs which internally call cuda runtime calls like below.
We are not looking to have any openmp runtime dependency on flang-rt APIs as in #145465

"""
void RTDEF(CUFDescriptorSync)(Descriptor *dst, const Descriptor *src,
const char *sourceFile, int sourceLine) {
std::size_t count{src->SizeInBytes()};
CUDA_REPORT_IF_ERROR(cudaMemcpy(
(void *)dst, (const void *)src, count, cudaMemcpyHostToDevice));
}
"""

@clementval
Copy link
Contributor

I don't think that adding OpenMP dependencies in flang runtime is a good idea.

By the way the box address can be retrieved without runtime call.

}
}

void *RTDEF(DescriptorGetBaseAddress)(
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

As Valentin said, there is fir.box_addr operation that allows taking the base address from a descriptor.

return baseAddr;
}

std::size_t RTDEF(DescriptorGetDataSizeInBytes)(
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is fir.box_total_elements (probably not implemented end-to-end right now) and fir.box_elesize, which can be used to compute the data size in bytes (of course, assuming that the data is contiguous).

So there are existing operations that should allow you to get all the data for omp_target_memcpy invocation and insert it in the compiler generated code rather than doing it in the Fortran runtime.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks @clementval and @vzakhari for feedback. Will check if these fir ops work in our scenario e2e. Closing the pull request for now.

@skc7 skc7 closed this Aug 20, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
flang:fir-hlfir flang Flang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

5 participants