Skip to content

[MLIR][OpenMP] Add Lowering support for OpenMP Declare Mapper directive #117046

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 10 commits into from
Feb 18, 2025

Conversation

TIFitis
Copy link
Member

@TIFitis TIFitis commented Nov 20, 2024

This patch adds HLFIR/FIR lowering support for OpenMP Declare Mapper directive.
Depends on #117045.

@llvmbot
Copy link
Member

llvmbot commented Nov 20, 2024

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

Author: Akash Banerjee (TIFitis)

Changes

This patch adds HLFIR/FIR lowering support for OpenMP Declare Mapper directive.


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

4 Files Affected:

  • (modified) flang/lib/Lower/OpenMP/OpenMP.cpp (+34-1)
  • (modified) flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp (+2-1)
  • (modified) flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90 (+4-4)
  • (added) flang/test/Lower/OpenMP/declare-mapper.f90 (+31)
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index a2779213a1a15a..c33bf0c9ea7a08 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -39,6 +39,7 @@
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
+#include <string>
 
 using namespace Fortran::lower::omp;
 
@@ -2701,7 +2702,39 @@ static void
 genOMP(lower::AbstractConverter &converter, lower::SymMap &symTable,
        semantics::SemanticsContext &semaCtx, lower::pft::Evaluation &eval,
        const parser::OpenMPDeclareMapperConstruct &declareMapperConstruct) {
-  TODO(converter.getCurrentLocation(), "OpenMPDeclareMapperConstruct");
+  fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
+  lower::StatementContext stmtCtx;
+  const auto &spec =
+      std::get<parser::OmpDeclareMapperSpecifier>(declareMapperConstruct.t);
+  const auto &mapperName{std::get<std::optional<parser::Name>>(spec.t)};
+  const auto &varType{std::get<parser::TypeSpec>(spec.t)};
+  const auto &varName{std::get<parser::Name>(spec.t)};
+  std::stringstream mapperNameStr;
+  if (mapperName.has_value()) {
+    mapperNameStr << mapperName->ToString();
+  } else {
+    mapperNameStr << "default_"
+                  << varType.declTypeSpec->derivedTypeSpec().name().ToString();
+  }
+
+  mlir::OpBuilder::InsertPoint insPt = firOpBuilder.saveInsertionPoint();
+  firOpBuilder.setInsertionPointToStart(converter.getModuleOp().getBody());
+  auto mlirType = converter.genType(varType.declTypeSpec->derivedTypeSpec());
+  auto varVal = firOpBuilder.createTemporaryAlloc(
+      converter.getCurrentLocation(), mlirType, varName.ToString());
+  symTable.addSymbol(*varName.symbol, varVal);
+
+  mlir::omp::DeclareMapperOperands clauseOps;
+  const auto *clauseList{
+      parser::Unwrap<parser::OmpClauseList>(declareMapperConstruct.t)};
+  List<Clause> clauses = makeClauses(*clauseList, semaCtx);
+  ClauseProcessor cp(converter, semaCtx, clauses);
+  cp.processMap(converter.getCurrentLocation(), stmtCtx, clauseOps);
+  auto declMapperOp = firOpBuilder.create<mlir::omp::DeclareMapperOp>(
+      converter.getCurrentLocation(), mapperNameStr.str(), varVal, mlirType,
+      clauseOps.mapVars);
+  converter.getMLIRSymbolTable()->insert(declMapperOp.getOperation());
+  firOpBuilder.restoreInsertionPoint(insPt);
 }
 
 static void
diff --git a/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp b/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp
index 4575c90e34acdd..01ffb40daa4aa2 100644
--- a/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp
+++ b/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp
@@ -447,7 +447,8 @@ class MapInfoFinalizationPass
     for (auto *user : mapOp->getUsers()) {
       if (llvm::isa<mlir::omp::TargetOp, mlir::omp::TargetDataOp,
                     mlir::omp::TargetUpdateOp, mlir::omp::TargetExitDataOp,
-                    mlir::omp::TargetEnterDataOp>(user))
+                    mlir::omp::TargetEnterDataOp, mlir::omp::DeclareMapperOp>(
+              user))
         return user;
 
       if (auto mapUser = llvm::dyn_cast<mlir::omp::MapInfoOp>(user))
diff --git a/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90 b/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90
index 5ae48ff7360482..13a4da5849f8c0 100644
--- a/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90
+++ b/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90
@@ -10,7 +10,7 @@ subroutine declare_mapper_1
  type my_type
    integer              :: num_vals
    integer, allocatable :: values(:)
- end type 
+ end type
 
  type my_type2
    type (my_type)        :: my_type_var
@@ -21,7 +21,7 @@ subroutine declare_mapper_1
   type (my_type2)        :: t
   real                   :: x, y(nvals)
   !$omp declare mapper (my_type :: var) map (var, var%values (1:var%num_vals))
-!CHECK: not yet implemented: OpenMPDeclareMapperConstruct
+!CHECK: not yet implemented: lowering symbol to HLFIR
 end subroutine declare_mapper_1
 
 
@@ -31,7 +31,7 @@ subroutine declare_mapper_2
  type my_type
    integer              :: num_vals
    integer, allocatable :: values(:)
- end type 
+ end type
 
  type my_type2
    type (my_type)        :: my_type_var
@@ -43,5 +43,5 @@ subroutine declare_mapper_2
   real                   :: x, y(nvals)
   !$omp declare mapper (my_mapper : my_type2 :: v) map (v%arr, x, y(:)) &
   !$omp&                map (alloc : v%temp)
-!CHECK: not yet implemented: OpenMPDeclareMapperConstruct
+!CHECK: not yet implemented: lowering symbol to HLFIR
 end subroutine declare_mapper_2
diff --git a/flang/test/Lower/OpenMP/declare-mapper.f90 b/flang/test/Lower/OpenMP/declare-mapper.f90
new file mode 100644
index 00000000000000..fd018b4fbb0e0f
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-mapper.f90
@@ -0,0 +1,31 @@
+! This test checks lowering of OpenMP declare mapper Directive.
+
+! RUN: split-file %s %t
+! RUN: %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=50 %t/declare-mapper-1.f90 -o - | FileCheck %t/declare-mapper-1.f90
+! RUN  %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=50 %t/declare-mapper-2.f90 -o - | FileCheck %t/declare-mapper-2.f90
+
+!--- declare-mapper-1.f90
+subroutine mapper
+   implicit none
+   type my_type
+      integer, pointer :: my_buffer
+      integer :: my_buffer_size
+   end type
+   !CHECK: %[[MY_VAR:.*]] = fir.alloca ![[VAR_TYPE:.*]] {bindc_name = "my_var"}
+   !CHECK: %[[MAP_INFO:.*]] = omp.map.info var_ptr(%[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]>, ![[VAR_TYPE]]) map_clauses(tofrom) capture(ByRef) -> !fir.ref<![[VAR_TYPE]]> {name = "my_var"}
+   !CHECK: omp.declare_mapper @my_mapper : %[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]> : ![[VAR_TYPE]] map_entries(%[[MAP_INFO]] : !fir.ref<![[VAR_TYPE]]>)
+   !$omp DECLARE MAPPER(my_mapper : my_type :: my_var) map(tofrom : my_var)
+end subroutine
+
+!--- declare-mapper-2.f90
+subroutine mapper_default
+   implicit none
+   type my_type
+      integer, pointer :: my_buffer
+      integer :: my_buffer_size
+   end type
+   !CHECK: %[[MY_VAR:.*]] = fir.alloca ![[VAR_TYPE:.*]] {bindc_name = "my_var"}
+   !CHECK: %[[MAP_INFO:.*]] = omp.map.info var_ptr(%[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]>, ![[VAR_TYPE]]) map_clauses(tofrom) capture(ByRef) -> !fir.ref<![[VAR_TYPE]]> {name = "my_var"}
+   !CHECK: omp.declare_mapper @default_my_type : %[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]> : ![[VAR_TYPE]] map_entries(%[[MAP_INFO]] : !fir.ref<![[VAR_TYPE]]>)
+   !$omp DECLARE MAPPER(my_type :: my_var) map(tofrom : my_var)
+end subroutine
\ No newline at end of file

@llvmbot
Copy link
Member

llvmbot commented Nov 20, 2024

@llvm/pr-subscribers-flang-openmp

Author: Akash Banerjee (TIFitis)

Changes

This patch adds HLFIR/FIR lowering support for OpenMP Declare Mapper directive.


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

4 Files Affected:

  • (modified) flang/lib/Lower/OpenMP/OpenMP.cpp (+34-1)
  • (modified) flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp (+2-1)
  • (modified) flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90 (+4-4)
  • (added) flang/test/Lower/OpenMP/declare-mapper.f90 (+31)
diff --git a/flang/lib/Lower/OpenMP/OpenMP.cpp b/flang/lib/Lower/OpenMP/OpenMP.cpp
index a2779213a1a15a..c33bf0c9ea7a08 100644
--- a/flang/lib/Lower/OpenMP/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP/OpenMP.cpp
@@ -39,6 +39,7 @@
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
+#include <string>
 
 using namespace Fortran::lower::omp;
 
@@ -2701,7 +2702,39 @@ static void
 genOMP(lower::AbstractConverter &converter, lower::SymMap &symTable,
        semantics::SemanticsContext &semaCtx, lower::pft::Evaluation &eval,
        const parser::OpenMPDeclareMapperConstruct &declareMapperConstruct) {
-  TODO(converter.getCurrentLocation(), "OpenMPDeclareMapperConstruct");
+  fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
+  lower::StatementContext stmtCtx;
+  const auto &spec =
+      std::get<parser::OmpDeclareMapperSpecifier>(declareMapperConstruct.t);
+  const auto &mapperName{std::get<std::optional<parser::Name>>(spec.t)};
+  const auto &varType{std::get<parser::TypeSpec>(spec.t)};
+  const auto &varName{std::get<parser::Name>(spec.t)};
+  std::stringstream mapperNameStr;
+  if (mapperName.has_value()) {
+    mapperNameStr << mapperName->ToString();
+  } else {
+    mapperNameStr << "default_"
+                  << varType.declTypeSpec->derivedTypeSpec().name().ToString();
+  }
+
+  mlir::OpBuilder::InsertPoint insPt = firOpBuilder.saveInsertionPoint();
+  firOpBuilder.setInsertionPointToStart(converter.getModuleOp().getBody());
+  auto mlirType = converter.genType(varType.declTypeSpec->derivedTypeSpec());
+  auto varVal = firOpBuilder.createTemporaryAlloc(
+      converter.getCurrentLocation(), mlirType, varName.ToString());
+  symTable.addSymbol(*varName.symbol, varVal);
+
+  mlir::omp::DeclareMapperOperands clauseOps;
+  const auto *clauseList{
+      parser::Unwrap<parser::OmpClauseList>(declareMapperConstruct.t)};
+  List<Clause> clauses = makeClauses(*clauseList, semaCtx);
+  ClauseProcessor cp(converter, semaCtx, clauses);
+  cp.processMap(converter.getCurrentLocation(), stmtCtx, clauseOps);
+  auto declMapperOp = firOpBuilder.create<mlir::omp::DeclareMapperOp>(
+      converter.getCurrentLocation(), mapperNameStr.str(), varVal, mlirType,
+      clauseOps.mapVars);
+  converter.getMLIRSymbolTable()->insert(declMapperOp.getOperation());
+  firOpBuilder.restoreInsertionPoint(insPt);
 }
 
 static void
diff --git a/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp b/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp
index 4575c90e34acdd..01ffb40daa4aa2 100644
--- a/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp
+++ b/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp
@@ -447,7 +447,8 @@ class MapInfoFinalizationPass
     for (auto *user : mapOp->getUsers()) {
       if (llvm::isa<mlir::omp::TargetOp, mlir::omp::TargetDataOp,
                     mlir::omp::TargetUpdateOp, mlir::omp::TargetExitDataOp,
-                    mlir::omp::TargetEnterDataOp>(user))
+                    mlir::omp::TargetEnterDataOp, mlir::omp::DeclareMapperOp>(
+              user))
         return user;
 
       if (auto mapUser = llvm::dyn_cast<mlir::omp::MapInfoOp>(user))
diff --git a/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90 b/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90
index 5ae48ff7360482..13a4da5849f8c0 100644
--- a/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90
+++ b/flang/test/Lower/OpenMP/Todo/omp-declare-mapper.f90
@@ -10,7 +10,7 @@ subroutine declare_mapper_1
  type my_type
    integer              :: num_vals
    integer, allocatable :: values(:)
- end type 
+ end type
 
  type my_type2
    type (my_type)        :: my_type_var
@@ -21,7 +21,7 @@ subroutine declare_mapper_1
   type (my_type2)        :: t
   real                   :: x, y(nvals)
   !$omp declare mapper (my_type :: var) map (var, var%values (1:var%num_vals))
-!CHECK: not yet implemented: OpenMPDeclareMapperConstruct
+!CHECK: not yet implemented: lowering symbol to HLFIR
 end subroutine declare_mapper_1
 
 
@@ -31,7 +31,7 @@ subroutine declare_mapper_2
  type my_type
    integer              :: num_vals
    integer, allocatable :: values(:)
- end type 
+ end type
 
  type my_type2
    type (my_type)        :: my_type_var
@@ -43,5 +43,5 @@ subroutine declare_mapper_2
   real                   :: x, y(nvals)
   !$omp declare mapper (my_mapper : my_type2 :: v) map (v%arr, x, y(:)) &
   !$omp&                map (alloc : v%temp)
-!CHECK: not yet implemented: OpenMPDeclareMapperConstruct
+!CHECK: not yet implemented: lowering symbol to HLFIR
 end subroutine declare_mapper_2
diff --git a/flang/test/Lower/OpenMP/declare-mapper.f90 b/flang/test/Lower/OpenMP/declare-mapper.f90
new file mode 100644
index 00000000000000..fd018b4fbb0e0f
--- /dev/null
+++ b/flang/test/Lower/OpenMP/declare-mapper.f90
@@ -0,0 +1,31 @@
+! This test checks lowering of OpenMP declare mapper Directive.
+
+! RUN: split-file %s %t
+! RUN: %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=50 %t/declare-mapper-1.f90 -o - | FileCheck %t/declare-mapper-1.f90
+! RUN  %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=50 %t/declare-mapper-2.f90 -o - | FileCheck %t/declare-mapper-2.f90
+
+!--- declare-mapper-1.f90
+subroutine mapper
+   implicit none
+   type my_type
+      integer, pointer :: my_buffer
+      integer :: my_buffer_size
+   end type
+   !CHECK: %[[MY_VAR:.*]] = fir.alloca ![[VAR_TYPE:.*]] {bindc_name = "my_var"}
+   !CHECK: %[[MAP_INFO:.*]] = omp.map.info var_ptr(%[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]>, ![[VAR_TYPE]]) map_clauses(tofrom) capture(ByRef) -> !fir.ref<![[VAR_TYPE]]> {name = "my_var"}
+   !CHECK: omp.declare_mapper @my_mapper : %[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]> : ![[VAR_TYPE]] map_entries(%[[MAP_INFO]] : !fir.ref<![[VAR_TYPE]]>)
+   !$omp DECLARE MAPPER(my_mapper : my_type :: my_var) map(tofrom : my_var)
+end subroutine
+
+!--- declare-mapper-2.f90
+subroutine mapper_default
+   implicit none
+   type my_type
+      integer, pointer :: my_buffer
+      integer :: my_buffer_size
+   end type
+   !CHECK: %[[MY_VAR:.*]] = fir.alloca ![[VAR_TYPE:.*]] {bindc_name = "my_var"}
+   !CHECK: %[[MAP_INFO:.*]] = omp.map.info var_ptr(%[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]>, ![[VAR_TYPE]]) map_clauses(tofrom) capture(ByRef) -> !fir.ref<![[VAR_TYPE]]> {name = "my_var"}
+   !CHECK: omp.declare_mapper @default_my_type : %[[MY_VAR]] : !fir.ref<![[VAR_TYPE]]> : ![[VAR_TYPE]] map_entries(%[[MAP_INFO]] : !fir.ref<![[VAR_TYPE]]>)
+   !$omp DECLARE MAPPER(my_type :: my_var) map(tofrom : my_var)
+end subroutine
\ No newline at end of file

Comment on lines 2725 to 2726
auto varVal = firOpBuilder.createTemporaryAlloc(
converter.getCurrentLocation(), mlirType, varName.ToString());
Copy link
Contributor

Choose a reason for hiding this comment

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

Sorry I didn't notice this before.

So far as I understand, this will create the fir.alloca and hlfir.declare at the beginning of the MLIR module, not nested in any intermediate operation.

How do you intend to lower this to LLVMIR? We would normally nest these in some kind of "function-like" wrapper operation e.g. func.func fir.global omp.private etc. I wonder if the declare mapper operation needs a nested region for this allocation (like we do for omp.private and omp.declare_reduction).

Copy link
Member Author

Choose a reason for hiding this comment

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

From what I understood when trying to implement this lowering:
The declMapperOp's parentOp must have a symbol table as it declares a new symbol, this is why I hoisted it to the ModuleOp from the funcOp.

The declMapperOp also relies on a var definition which is immediately used in it's map clause. I can't put this var inside a region, as then it can't be used in the map clause. As such, this new alloc is also moved to just before the declMapperOp.

Do you have any suggestion to do this in a different way?

Copy link
Contributor

@tblah tblah Nov 25, 2024

Choose a reason for hiding this comment

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

Unfortunately I am not very familiar with the target offload parts of OpenMP so I might have missunderstood something here.

OpenMP 5.2 (section 5.8.8) says

The visibility and accessibility of this declaration are the same as those of a variable declared at the same location in the program.

Therefore I don't think it is right to be pushing this to the module scope in case there are conflicting mappers declared in different scopes in the same module.

If I understand correctly, the declare mapper directive declares var:

The variable declared by var is available for use in all map clauses on the directive, and no part of the variable to be mapped is mapped by default.

So it should be treated as though this parse node is a variable declaration so createTemporaryAlloc would be appropriate if this is inside of a function, but if it is not, this should probably be a fir.global (like if you declared a module- or file-level variable).

Copy link
Member Author

Choose a reason for hiding this comment

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

Therefore I don't think it is right to be pushing this to the module scope in case there are conflicting mappers declared in different scopes in the same module.

It isn't my intention to move the declMapperOp to the module scope, but rather MLIR is forcing me to. If there's an alternative I'm missing where I can keep the declMapperOp within the func/block where it was declared, I'd much rather do that. In case there is no way to do that in MLIR, I can prepend the name of the function to the module scope declaration to prevent conflicts.

So it should be treated as though this parse node is a variable declaration so createTemporaryAlloc would be appropriate if this is inside of a function, but if it is not, this should probably be a fir.global (like if you declared a module- or file-level variable).

If we go forward with keeping the declMapperOp in the module scope, then I'll alter the var declaration to be a fir.global.

Copy link
Contributor

Choose a reason for hiding this comment

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

Isn't declare mapper providing a specification which is to be used for mapping in the associated regions? The variable in the declare mapper has its scope inside the construct since it is only used for specifying how to map. Similar to what @tblah is saying, you could do something like the following.

Assuming %x is the variable that is being mapped by the rules of the declare mapper:

omp.declare_mapper @default_my_type {
^bb0(%[[VAR:.*]]: [[TYPE]]):
  %v1 = omp.bounds
  %v2 = omp.map.info var_ptr(%[[VAR:.*]])
  omp.yield %v2
}

%x_map = omp.map.info var_ptr(%x) decl_map(@default_my_type)
omp.target_enter_data   map_entries(%x_map)

or

omp.declare_mapper @default_my_type {
^bb0(%[[VAR:.*]]: [[TYPE]]):
  %v1 = omp.bounds
  %v2 = omp.map.info var_ptr(%[[VAR:.*]])
  omp.yield %v2
}
omp.target_enter_data   map_entries(%x : @default_my_type)

The alternative would be to not have a declare mapper construct and apply the declare mapper during lowering itself. But this would be early lowering and we wont be able to share code with Clang as well.

Copy link
Contributor

@agozillon agozillon Nov 27, 2024

Choose a reason for hiding this comment

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

1. If a DECLARE MAPPER directive is not specified for a type DT, a predefined mapper exists for type DT as if the type DT had appeared in the directive as follows:

!$OMP DECLARE MAPPER (DT :: var) MAP (TOFROM: var)
2. If a variable is not a scalar then it is treated as if it had appeared in a map clause with a map-type of tofrom. Which is effectively equivalent to the following and extending declare mapper for non-derived types:
!$OMP DECLARE MAPPER (T :: var) MAP (TOFROM: var)

I think the keyword here is likely "as if", so as long as the effects are as described it's reasonable would be my (albeit terrible) reading, and if we really wanted to be exact about the wording we'd generate/embed our own equivalent pragmas to the above for all default mappings, and then lower them, so not just at the MLIR level. However, saying that I am not against defining a default declare mapper for all cases once it's in place, it might tidy things up a bit, but it may also be more complicated/trouble than it's worth, in either case I am fine with the approach of defining default declare mappers if we'd like to go down that route :-)

I'd also love it if whatever implementation we landed on was compatible with the OpenACC implementations documentation/approach to mapping descriptors via runtime calls, as I'd like to move towards that eventually when I have some time to dig into it and see if it's viable for us. I imagine it will be, I just don't know a ton about the region'd approach so hope it wouldn't be prohibitive of this.

Copy link
Member Author

Choose a reason for hiding this comment

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

I have made the required changes to both PRs to update the representation for the DeclMapperOp.

Some changes I've incorporated are:

Inside DeclMapperOp's region I've introduce a new DeclMapperInfoOp to which I've attached the MapClause. Not having any MapClause explicitly associated seemed weird to me, also walking through the region collecting all the MapInfoOps for processing in various places in the code base seemed like a bad design to me.

Also, instead of using the block arg, I've created a temporary alloca inside the region. This is to save the hassle of binding the block_arg to the symbol, binding the alloca is much more straightforward.

We can drop the entire DeclMapperOp including the region once it reaches OpenMPTOLLVMIRTranslation.

Let me what are your thoughts on this approach :)

Copy link
Member Author

Choose a reason for hiding this comment

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

@kiranchandramohan @tblah @skatrak @agozillon Any thoughts on this approach?

Copy link
Contributor

Choose a reason for hiding this comment

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

You should feel free to proceed with the approval of @agozillon and @skatrak. You can set up a quick call with them and describe your approach and come to a conclusion. You can report the outcomes somewhere in this patch or write in discourse.

Inside DeclMapperOp's region I've introduce a new DeclMapperInfoOp to which I've attached the MapClause. Not having any MapClause explicitly associated seemed weird to me, also walking through the region collecting all the MapInfoOps for processing in various places in the code base seemed like a bad design to me.

The idea would generally be to inline the whole declare mapper operation region, replacing the block_arg with the real variable that is going to be mapped. Would an alloca always be created for all kinds of mappings? If not committing to an alloca might mean you have to delete it in some circumstances.

We can drop the entire DeclMapperOp including the region once it reaches OpenMPTOLLVMIRTranslation.

As in just drop it without using it? Or using it create the @.omp_mapper._ZTS1T.deep_copy function for the declare mapper (#pragma omp declare mapper(deep_copy : T abc) map(abc, abc.ptr[ : abc.buf_size])) in the C example you gave below.

You have not talked about the following two points.

-> where we create the map_entries for the relevant operations (like target) for which the declare mapper implicitly applies. Currently, this is done during lowering. But in this patch you have solely focused on creating the declare mapper.
-> where the composition of map-types (map-type decay) from the map clause of declare mapper and the map clause of the relevant operation (like target) happens

Couple of other things that came to my mind :
-> Since declare mappers are in the specification section, it can also occur in modules. We have not added code to propagate it to the module file. If this is not urgent for you, I can fix it sometime.
-> You should test the case where the declare mapper is in the host subroutine

subroutine A
type t
end type
declare mapper(t)
contains
subroutine B
!$omp target 
! use a var of type t
!$omp end target
end subroutine

Copy link
Member Author

Choose a reason for hiding this comment

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

You should feel free to proceed with the approval of @agozillon and @skatrak.
Thanks for the go ahead, I'll consult with them before proceeding, and post any relevant findings here.

Inside DeclMapperOp's region I've introduce a new DeclMapperInfoOp to which I've attached the MapClause. Not having any MapClause explicitly associated seemed weird to me, also walking through the region collecting all the MapInfoOps for processing in various places in the code base seemed like a bad design to me.

The idea would generally be to inline the whole declare mapper operation region, replacing the block_arg with the real variable that is going to be mapped. Would an alloca always be created for all kinds of mappings? If not committing to an alloca might mean you have to delete it in some circumstances.

In case of DeclareMapper the mapping is handled in a ad-hoc basis. There is no real variable to be replaced later. The variable simply exists as a dummy to represent the mapping information for other real variables where the mapping would be used. However, this mapping isn't applied directly to those variables either, rather the runtime is notified that a mapper function exists for these variables.

We can drop the entire DeclMapperOp including the region once it reaches OpenMPTOLLVMIRTranslation.

As in just drop it without using it? Or using it create the @.omp_mapper._ZTS1T.deep_copy function for the declare mapper (#pragma omp declare mapper(deep_copy : T abc) map(abc, abc.ptr[ : abc.buf_size])) in the C example you gave below.
Yes, I mean't each declareMapperOp would be resolved to a function like @.omp_mapper._ZTS1T.deep_copy. By dropping it, I mean't it's region won't be separately lowered.

You have not talked about the following two points.

-> where we create the map_entries for the relevant operations (like target) for which the declare mapper implicitly applies. Currently, this is done during lowering. But in this patch you have solely focused on creating the declare mapper.
-> where the composition of map-types (map-type decay) from the map clause of declare mapper and the map clause of the relevant operation (like target) happens

I'll address the implicit mapping either in the DeclMapper lowering support for the mapClause or a separate PR soon after that.

The composition of mapClause AFAIK is likely handled in the OMPIRBuilder. If that's not the case, I'll address it in a future patch.

Couple of other things that came to my mind : -> Since declare mappers are in the specification section, it can also occur in modules. We have not added code to propagate it to the module file. If this is not urgent for you, I can fix it sometime.

Please feel free to do this if you're already familiar with what needs to be done, as I'm not. Otherwise, I can add it to my TODO list :)

-> You should test the case where the declare mapper is in the host subroutine

subroutine A
type t
end type
declare mapper(t)
contains
subroutine B
!$omp target 
! use a var of type t
!$omp end target
end subroutine

Thanks, I'll test this in the mapClause mapper support PR once I have it ready.

Many thanks for the review.

@TIFitis TIFitis force-pushed the users/akash/mapper_lowering branch from a9f42ac to 17ce48e Compare November 28, 2024 21:18
@TIFitis
Copy link
Member Author

TIFitis commented Nov 28, 2024

As @agozillon requested, here's a sample of how Clang lowers declare mapper.

C Code:

typedef struct {
  int *ptr;
  int buf_size;
} T;
#pragma omp declare mapper(deep_copy : T abc) map(abc, abc.ptr[ : abc.buf_size])
int main() {
  T xyz;
#pragma omp target data map(mapper(deep_copy), tofrom : xyz)
  {
  }
  return 0;
}

LLVM IR:

@.offload_sizes = private unnamed_addr constant [1 x i64] [i64 16]
@.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 3]

; Function Attrs: noinline nounwind optnone uwtable
define dso_local i32 @main() #0 {
entry:
  %retval = alloca i32, align 4
  %xyz = alloca %struct.T, align 8
  %.offload_baseptrs = alloca [1 x ptr], align 8
  %.offload_ptrs = alloca [1 x ptr], align 8
  %.offload_mappers = alloca [1 x ptr], align 8
  store i32 0, ptr %retval, align 4
  %0 = getelementptr inbounds [1 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
  store ptr %xyz, ptr %0, align 8
  %1 = getelementptr inbounds [1 x ptr], ptr %.offload_ptrs, i32 0, i32 0
  store ptr %xyz, ptr %1, align 8
  %2 = getelementptr inbounds [1 x ptr], ptr %.offload_mappers, i64 0, i64 0
  store ptr @.omp_mapper._ZTS1T.deep_copy, ptr %2, align 8
  %3 = getelementptr inbounds [1 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
  %4 = getelementptr inbounds [1 x ptr], ptr %.offload_ptrs, i32 0, i32 0
  call void @__tgt_target_data_begin_mapper(ptr @1, i64 -1, i32 1, ptr %3, ptr %4, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr %.offload_mappers)
  %5 = getelementptr inbounds [1 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
  %6 = getelementptr inbounds [1 x ptr], ptr %.offload_ptrs, i32 0, i32 0
  call void @__tgt_target_data_end_mapper(ptr @1, i64 -1, i32 1, ptr %5, ptr %6, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr %.offload_mappers)
  ret i32 0
}

; Function Attrs: noinline nounwind uwtable
define internal void @.omp_mapper._ZTS1T.deep_copy(ptr noundef %0, ptr noundef %1, ptr noundef %2, i64 noundef %3, i64 noundef %4, ptr noundef %5) #1 {
entry:
  %.addr = alloca ptr, align 8
  %.addr1 = alloca ptr, align 8
  %.addr2 = alloca ptr, align 8
  %.addr3 = alloca i64, align 8
  %.addr4 = alloca i64, align 8
  %.addr5 = alloca ptr, align 8
  store ptr %0, ptr %.addr, align 8
  store ptr %1, ptr %.addr1, align 8
  store ptr %2, ptr %.addr2, align 8
  store i64 %3, ptr %.addr3, align 8
  store i64 %4, ptr %.addr4, align 8
  store ptr %5, ptr %.addr5, align 8
  %6 = load i64, ptr %.addr3, align 8
  %7 = load ptr, ptr %.addr, align 8
  %8 = load ptr, ptr %.addr1, align 8
  %9 = load ptr, ptr %.addr2, align 8
  %10 = udiv exact i64 %6, 16
  %11 = getelementptr %struct.T, ptr %9, i64 %10
  %12 = load i64, ptr %.addr4, align 8
  %13 = load ptr, ptr %.addr5, align 8
  %omp.arrayinit.isarray = icmp sgt i64 %10, 1
  %14 = and i64 %12, 8
  %15 = icmp ne ptr %8, %9
  %16 = and i64 %12, 16
  %17 = icmp ne i64 %16, 0
  %18 = and i1 %15, %17
  %19 = or i1 %omp.arrayinit.isarray, %18
  %.omp.array..init..delete = icmp eq i64 %14, 0
  %20 = and i1 %19, %.omp.array..init..delete
  br i1 %20, label %.omp.array..init, label %omp.arraymap.head

.omp.array..init:                                 ; preds = %entry
  %21 = mul nuw i64 %10, 16
  %22 = and i64 %12, -4
  %23 = or i64 %22, 512
  call void @__tgt_push_mapper_component(ptr %7, ptr %8, ptr %9, i64 %21, i64 %23, ptr %13)
  br label %omp.arraymap.head

omp.arraymap.head:                                ; preds = %.omp.array..init, %entry
  %omp.arraymap.isempty = icmp eq ptr %9, %11
  br i1 %omp.arraymap.isempty, label %omp.done, label %omp.arraymap.body

omp.arraymap.body:                                ; preds = %omp.type.end19, %omp.arraymap.head
  %omp.arraymap.ptrcurrent = phi ptr [ %9, %omp.arraymap.head ], [ %omp.arraymap.next, %omp.type.end19 ]
  %ptr = getelementptr inbounds nuw %struct.T, ptr %omp.arraymap.ptrcurrent, i32 0, i32 0
  %ptr6 = getelementptr inbounds nuw %struct.T, ptr %omp.arraymap.ptrcurrent, i32 0, i32 0
  %24 = load ptr, ptr %ptr6, align 8
  %arrayidx = getelementptr inbounds nuw i32, ptr %24, i64 0
  %buf_size = getelementptr inbounds nuw %struct.T, ptr %omp.arraymap.ptrcurrent, i32 0, i32 1
  %25 = load i32, ptr %buf_size, align 8
  %conv = sext i32 %25 to i64
  %26 = mul nuw i64 %conv, 4
  %27 = getelementptr %struct.T, ptr %omp.arraymap.ptrcurrent, i32 1
  %28 = ptrtoint ptr %27 to i64
  %29 = ptrtoint ptr %omp.arraymap.ptrcurrent to i64
  %30 = sub i64 %28, %29
  %31 = sdiv exact i64 %30, ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
  %32 = call i64 @__tgt_mapper_num_components(ptr %7)
  %33 = shl i64 %32, 48
  %34 = add nuw i64 0, %33
  %35 = and i64 %12, 3
  %36 = icmp eq i64 %35, 0
  br i1 %36, label %omp.type.alloc, label %omp.type.alloc.else

omp.type.alloc:                                   ; preds = %omp.arraymap.body
  %37 = and i64 %34, -4
  br label %omp.type.end

omp.type.alloc.else:                              ; preds = %omp.arraymap.body
  %38 = icmp eq i64 %35, 1
  br i1 %38, label %omp.type.to, label %omp.type.to.else

omp.type.to:                                      ; preds = %omp.type.alloc.else
  %39 = and i64 %34, -3
  br label %omp.type.end

omp.type.to.else:                                 ; preds = %omp.type.alloc.else
  %40 = icmp eq i64 %35, 2
  br i1 %40, label %omp.type.from, label %omp.type.end

omp.type.from:                                    ; preds = %omp.type.to.else
  %41 = and i64 %34, -2
  br label %omp.type.end

omp.type.end:                                     ; preds = %omp.type.from, %omp.type.to.else, %omp.type.to, %omp.type.alloc
  %omp.maptype = phi i64 [ %37, %omp.type.alloc ], [ %39, %omp.type.to ], [ %41, %omp.type.from ], [ %34, %omp.type.to.else ]
  call void @__tgt_push_mapper_component(ptr %7, ptr %omp.arraymap.ptrcurrent, ptr %omp.arraymap.ptrcurrent, i64 %31, i64 %omp.maptype, ptr null)
  %42 = add nuw i64 281474976710659, %33
  %43 = and i64 %12, 3
  %44 = icmp eq i64 %43, 0
  br i1 %44, label %omp.type.alloc7, label %omp.type.alloc.else8

omp.type.alloc7:                                  ; preds = %omp.type.end
  %45 = and i64 %42, -4
  br label %omp.type.end12

omp.type.alloc.else8:                             ; preds = %omp.type.end
  %46 = icmp eq i64 %43, 1
  br i1 %46, label %omp.type.to9, label %omp.type.to.else10

omp.type.to9:                                     ; preds = %omp.type.alloc.else8
  %47 = and i64 %42, -3
  br label %omp.type.end12

omp.type.to.else10:                               ; preds = %omp.type.alloc.else8
  %48 = icmp eq i64 %43, 2
  br i1 %48, label %omp.type.from11, label %omp.type.end12

omp.type.from11:                                  ; preds = %omp.type.to.else10
  %49 = and i64 %42, -2
  br label %omp.type.end12

omp.type.end12:                                   ; preds = %omp.type.from11, %omp.type.to.else10, %omp.type.to9, %omp.type.alloc7
  %omp.maptype13 = phi i64 [ %45, %omp.type.alloc7 ], [ %47, %omp.type.to9 ], [ %49, %omp.type.from11 ], [ %42, %omp.type.to.else10 ]
  call void @__tgt_push_mapper_component(ptr %7, ptr %omp.arraymap.ptrcurrent, ptr %omp.arraymap.ptrcurrent, i64 16, i64 %omp.maptype13, ptr null)
  %50 = add nuw i64 281474976710675, %33
  %51 = and i64 %12, 3
  %52 = icmp eq i64 %51, 0
  br i1 %52, label %omp.type.alloc14, label %omp.type.alloc.else15

omp.type.alloc14:                                 ; preds = %omp.type.end12
  %53 = and i64 %50, -4
  br label %omp.type.end19

omp.type.alloc.else15:                            ; preds = %omp.type.end12
  %54 = icmp eq i64 %51, 1
  br i1 %54, label %omp.type.to16, label %omp.type.to.else17

omp.type.to16:                                    ; preds = %omp.type.alloc.else15
  %55 = and i64 %50, -3
  br label %omp.type.end19

omp.type.to.else17:                               ; preds = %omp.type.alloc.else15
  %56 = icmp eq i64 %51, 2
  br i1 %56, label %omp.type.from18, label %omp.type.end19

omp.type.from18:                                  ; preds = %omp.type.to.else17
  %57 = and i64 %50, -2
  br label %omp.type.end19

omp.type.end19:                                   ; preds = %omp.type.from18, %omp.type.to.else17, %omp.type.to16, %omp.type.alloc14
  %omp.maptype20 = phi i64 [ %53, %omp.type.alloc14 ], [ %55, %omp.type.to16 ], [ %57, %omp.type.from18 ], [ %50, %omp.type.to.else17 ]
  call void @__tgt_push_mapper_component(ptr %7, ptr %ptr, ptr %arrayidx, i64 %26, i64 %omp.maptype20, ptr null)
  %omp.arraymap.next = getelementptr %struct.T, ptr %omp.arraymap.ptrcurrent, i32 1
  %omp.arraymap.isdone = icmp eq ptr %omp.arraymap.next, %11
  br i1 %omp.arraymap.isdone, label %omp.arraymap.exit, label %omp.arraymap.body

omp.arraymap.exit:                                ; preds = %omp.type.end19
  %omp.arrayinit.isarray21 = icmp sgt i64 %10, 1
  %58 = and i64 %12, 8
  %.omp.array..del..delete = icmp ne i64 %58, 0
  %59 = and i1 %omp.arrayinit.isarray21, %.omp.array..del..delete
  br i1 %59, label %.omp.array..del, label %omp.done

.omp.array..del:                                  ; preds = %omp.arraymap.exit
  %60 = mul nuw i64 %10, 16
  %61 = and i64 %12, -4
  %62 = or i64 %61, 512
  call void @__tgt_push_mapper_component(ptr %7, ptr %8, ptr %9, i64 %60, i64 %62, ptr %13)
  br label %omp.done

omp.done:                                         ; preds = %.omp.array..del, %omp.arraymap.exit, %omp.arraymap.head
  ret void
}

@TIFitis
Copy link
Member Author

TIFitis commented Dec 10, 2024

@kiranchandramohan I discussed the current approach with @skatrak today. When trying to implement the mapper lowering for the map clause, it became apparent that we need to add the declMapperOp name to the SymbolTable. As such, we would also need to hoist the declareMapperOp to the ModuleOp.

I am however struggling with name mangling the declMapperOp such that scoping information is preserved, and we don't have clashing declMapperOps from different nested scopes with the same name.

Take the following example:

program my_prog
   type my_type
      integer               :: num
   end type
   !$omp declare mapper (my_mapper : my_type :: my_var) map (my_var)
contains
   subroutine test
      type(my_type)            :: xyz
      !$omp target enter data map(mapper(my_mapper), to: xyz)
   end subroutine test
end program my_prog

Here after mangling the, declMapperOp symbol name becomes _QQFmy_mapper. But when trying to mangle the name that occurs in the mapClause before lookup results in _QQFFtestmy_mapper.

Do you know any mechanism in Fortran lowering that could help resolve this issue?

@kiranchandramohan
Copy link
Contributor

kiranchandramohan commented Dec 11, 2024

@kiranchandramohan I discussed the current approach with @skatrak today. When trying to implement the mapper lowering for the map clause, it became apparent that we need to add the declMapperOp name to the SymbolTable. As such, we would also need to hoist the declareMapperOp to the ModuleOp.

I am however struggling with name mangling the declMapperOp such that scoping information is preserved, and we don't have clashing declMapperOps from different nested scopes with the same name.

Take the following example:

program my_prog
   type my_type
      integer               :: num
   end type
   !$omp declare mapper (my_mapper : my_type :: my_var) map (my_var)
contains
   subroutine test
      type(my_type)            :: xyz
      !$omp target enter data map(mapper(my_mapper), to: xyz)
   end subroutine test
end program my_prog

Here after mangling the, declMapperOp symbol name becomes _QQFmy_mapper. But when trying to mangle the name that occurs in the mapClause before lookup results in _QQFFtestmy_mapper.

Do you know any mechanism in Fortran lowering that could help resolve this issue?

The frontend symbols also are implementing a symbol table. For the example you showed, I am assuming the symbol for my_mapper in target enter data map is same as the symbol for my_mapper in declare mapper either directly or via host-association. So, naturally they should be mangling to the same name. This is also similar to a variable defined in the program being used in the contained subroutine. The mangling will follow the program's mangling and not the subroutine's mangling.

Are the symbols not being resolved properly? Or did I miss a point?

@skatrak
Copy link
Member

skatrak commented Dec 11, 2024

After the discussion with @TIFitis yesterday, I think the overall approach currently proposed makes sense. I would just like to share a few related thoughts not for this PR stack, but rather some longer term potential improvements that might make sense to do:

  • I think the new omp.declare_mapper_info operation currently helps not having to introduce too many new features at once, and rely on the existing infrastructure as much as possible, but I believe eventually this should be replaced with an omp.yield of the necessary omp.map.info operations and potentially some changes to the verification and handling of map clauses.
  • Long-term I was thinking that we may want to represent type-dependent handling of map clauses in a way more similar to what we do for privatization. Initially I was thinking that we could use the omp.declare_mapper operation for this, but it's true that it's only intended to be used for derived types / structures. However, my thinking is that if we introduced something like an omp.map operation holding region(s) that define mapping for a given type of variable, we may be able to avoid the need for fix-up passes to add new mapped variables for e.g. allocatables or other complex types. Ideally, it would be able to integrate with omp.declare_mapper operations and unify map-related handling so that there's fewer corner cases or at least centralize that kind of logic. This is only a high-level idea at this point, not something I have actually spent time thinking through in detail, but perhaps something worth exploring at some point.

@kparzysz
Copy link
Contributor

Are the symbols not being resolved properly? Or did I miss a point?

I'm working on something different at the moment, but for debugging purposes I changed the tree dumper to show symbol information too. Here's the output for this code. It shows that the variable my_mapper in the declare target, and in the map clause later on use the same symbol: 0x61f689c35c30.

Program -> ProgramUnit -> MainProgram
| ProgramStmt -> Name = '(0x61f689c354b0) [my_prog: MainProgram]'
| SpecificationPart
| | ImplicitPart ->
| | DeclarationConstruct -> SpecificationConstruct -> DerivedTypeDef
| | | DerivedTypeStmt
| | | | Name = '(0x61f689c35870) [my_type: DerivedType components: num]'
| | | ComponentDefStmt -> DataComponentDefStmt
| | | | DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec ->
| | | | ComponentOrFill -> ComponentDecl
| | | | | Name = '(0x61f689c35a50) [num size=4 offset=0: ObjectEntity type: INTEGER(4)]'
| | | EndTypeStmt ->
| | DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareMapperConstruct
| | | Verbatim
| | | OmpDeclareMapperSpecifier
| | | | Name = '(0x61f689c35c30) [my_mapper: Misc ConstructName]'
| | | | TypeSpec -> DerivedTypeSpec
| | | | | Name = '(0x61f689c35870) [my_type: DerivedType components: num]'
| | | | Name = '(0x61f689c35e10) [my_var (OmpMapToFrom) size=4 offset=0: ObjectEntity type: TYPE(my_type)]'
| | | OmpClauseList -> OmpClause -> Map -> OmpMapClause
| | | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = '(0x61f689c35e10) [my_var (OmpMapToFrom) size=4 offset=0: ObjectEntity type: TYPE(my_type)]'
| | | | bool = 'true'
| ExecutionPart -> Block
| InternalSubprogramPart
| | ContainsStmt
| | InternalSubprogram -> SubroutineSubprogram
| | | SubroutineStmt
| | | | Name = '(0x61f689c35690) [test (Subroutine): Subprogram ()]'
| | | SpecificationPart
| | | | ImplicitPart ->
| | | | DeclarationConstruct -> SpecificationConstruct -> TypeDeclarationStmt
| | | | | DeclarationTypeSpec -> Type
| | | | | | DerivedTypeSpec
| | | | | | | Name = '(0x61f689c35870) [my_type: DerivedType components: num]'
| | | | | EntityDecl
| | | | | | Name = '(0x61f689c361d0) [xyz (OmpMapTo) size=4 offset=0: ObjectEntity type: TYPE(my_type)]'
| | | ExecutionPart -> Block
| | | | ExecutionPartConstruct -> ExecutableConstruct -> OpenMPConstruct -> OpenMPStandaloneConstruct -> OpenMPSimpleStandaloneConstruct
| | | | | OmpSimpleStandaloneDirective -> llvm::omp::Directive = target enter data
| | | | | OmpClauseList -> OmpClause -> Map -> OmpMapClause
| | | | | | Modifier -> OmpMapper -> Name = '(0x61f689c35c30) [my_mapper: Misc ConstructName]'
| | | | | | Modifier -> OmpMapType -> Value = To
| | | | | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = '(0x61f689c361d0) [xyz (OmpMapTo) size=4 offset=0: ObjectEntity type: TYPE(my_type)]'
| | | | | | bool = 'true'
| | | EndSubroutineStmt -> Name = 'test'
| EndProgramStmt -> Name = 'my_prog'

@TIFitis
Copy link
Member Author

TIFitis commented Dec 11, 2024

@kiranchandramohan @kparzysz I guess I must be using the name mangler in an incorrect way then. I've added the code snippets I am using when lowering and later performing a lookup. Please let me know what would be the correct way of doing this.

When lowering DeclMapperOp:

static void
genOMP(lower::AbstractConverter &converter, lower::SymMap &symTable,
       semantics::SemanticsContext &semaCtx, lower::pft::Evaluation &eval,
       const parser::OpenMPDeclareMapperConstruct &declareMapperConstruct) {
  fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
  const auto &spec =
      std::get<parser::OmpDeclareMapperSpecifier>(declareMapperConstruct.t);
  const auto &mapperName{std::get<std::optional<parser::Name>>(spec.t)};
  std::string mapperNameStr;
  if (mapperName.has_value())
    mapperNameStr = mapperName->ToString();
  else
    mapperNameStr =
        "default_" + varType.declTypeSpec->derivedTypeSpec().name().ToString();

  // This is my current implementation.
  // It returns _QQFmy_mapper as mapperNameStr.
  mapperNameStr = converter.mangleName(*mapperName->symbol);

  auto declMapperOp = firOpBuilder.create<mlir::omp::DeclareMapperOp>(
      loc, mapperNameStr, mlirType);
  converter.getMLIRSymbolTable()->insert(declMapperOp);

  // This code produces the error:
  // LLVM ERROR: /home/akash/Documents/llvm-project/flang/lib/Lower/Mangler.cpp:184: not yet implemented: symbol mangling
  // mapperNameStr = converter.mangleName(*mapperName->symbol);
  
  // This code produces the error:
  // ** symbol not properly mapped **
  // symTable.lookupSymbol(mapperName->symbol).dump();

When trying to lookup DeclareMapper from the mapClause mapper:

      auto mapperIdName = mappers->front().v.id().symbol->name().ToString();
      mapperIdName = converter.mangleName(mapperIdName);

      // Here mapperIdName returns _QQFFtestmy_mapper which fails the
      // lookup inside the asser.

      // assert(converter.getMLIRSymbolTable()->lookup(mapperIdName) &&
      //        "mapper not found");
      mapperId = mlir::FlatSymbolRefAttr::get(&converter.getMLIRContext(),
                                              mapperIdName);

@kparzysz
Copy link
Contributor

  1. If you want to create an MLIR op for the mapper, you can give it some name (you can mangle it if you want). You will need to store it in the MLIR somewhere/somehow.
  2. If you want to record the use of a mapper in a clause, you can just use the name of the mapper in the same form as in (1), i.e. you decide how to generate the name here.
  3. If you want to actually apply the mapper, you need to find it, but it's up to you how you do it: if you're looking for it in the MLIR module then your code must have put it there, if you're looking for it in the AST then you can find the declaration of the mapper by the symbol associated with its name.

@kparzysz
Copy link
Contributor

I looked at the clang code and it's not clear to me why they do it at runtime. Mappers can refer to other mappers (for sub-objects), but all mappers and all type layouts should be present in (or obtainable from) the AST[1]. In other words, clang should be able to emit all the "expanded" map clauses (i.e. after application of the mappers) at compile-time. It's possible that instead of emitting all that code inline, they put it in a runtime function to save code size.

[1] The mapper definition must be visible at the time of use in a clause, so all mappers (recursively) should be visible.

@TIFitis
Copy link
Member Author

TIFitis commented Jan 28, 2025

@kiranchandramohan @skatrak @agozillon @kparzysz Hi everyone! I was on vacation these last couple of weeks, apologies for the delay.

I have added PR #124746 with which we have all the implementation in place for declare mappers.

It would be great to resume the review process for this PR stack.

Thanks :)

Note: I'll add a couple of patches to support the implicit default mapping support for declare mappers.

@TIFitis
Copy link
Member Author

TIFitis commented Feb 3, 2025

Polite request for review 🙂

Copy link
Contributor

@tblah tblah left a comment

Choose a reason for hiding this comment

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

I have some minor suggestions on the code. Please wait for review from somebody with more familiarity with omp target things, and this is conditional on the design of the MLIR operation being approved.

Copy link
Member

@skatrak skatrak left a comment

Choose a reason for hiding this comment

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

LGTM, thank you.

TIFitis added a commit that referenced this pull request Feb 18, 2025
This patch adds the OMP.DeclareMapperOp to MLIR.
The HLFIR/FIR lowering for Declare Mapper is available here #117046.
Base automatically changed from users/akash/mapper_op to main February 18, 2025 16:11
@TIFitis TIFitis force-pushed the users/akash/mapper_lowering branch from aa7f4aa to afb17c9 Compare February 18, 2025 16:22
@TIFitis TIFitis merged commit 9905728 into main Feb 18, 2025
5 of 7 checks passed
@TIFitis TIFitis deleted the users/akash/mapper_lowering branch February 18, 2025 16:36
TIFitis added a commit that referenced this pull request Feb 18, 2025
This patch adds the mapper field to the omp.map.info op.

Depends on #117046.
wldfngrs pushed a commit to wldfngrs/llvm-project that referenced this pull request Feb 19, 2025
This patch adds the OMP.DeclareMapperOp to MLIR.
The HLFIR/FIR lowering for Declare Mapper is available here llvm#117046.
wldfngrs pushed a commit to wldfngrs/llvm-project that referenced this pull request Feb 19, 2025
…ve (llvm#117046)

This patch adds HLFIR/FIR lowering support for OpenMP Declare Mapper
directive.
Depends on llvm#117045.
wldfngrs pushed a commit to wldfngrs/llvm-project that referenced this pull request Feb 19, 2025
This patch adds the mapper field to the omp.map.info op.

Depends on llvm#117046.
TIFitis added a commit to TIFitis/llvm-project that referenced this pull request Feb 20, 2025
This patch adds the mapper field to the omp.map.info op.

Depends on llvm#117046.
searlmc1 pushed a commit to ROCm/llvm-project that referenced this pull request Feb 22, 2025
This patch adds the mapper field to the omp.map.info op.

Depends on llvm#117046.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
flang:fir-hlfir flang:openmp flang Flang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

7 participants