[llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Previous Topic Next Topic
 
classic Classic list List threaded Threaded
11 messages Options
Reply | Threaded
Open this post in threaded view
|

[llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev

RFC: Representing the target device information in the LLVM IR

===========================================================================

 

Why this RFC change?

=================

The target device information needs to be passed to the LLVM backend when OpenMP backend outlining is enabled. For example, for multiple target devices, the target compilation has to generate a single host to support all the targets. In order to make sure all the target outlined functions have the same interface, the information of all the target architectures is needed during host and target compilation. In the following example, the firstprivate variable ‘d’ is passed by value under x86_64-mic and passed by reference under i386-pc-linux-gnu. In order to avoid this inconsistency, the compiler needs all the target architecture information so it can find a common interface. In this example, it will change the x86_64-mic interface for ‘d’ to pass by reference.

 

Existing code: 64-bit firstprivate variable

 

void foo() {

 double d = 1.0;

 #pragma omp target firstprivate(d)

 {}

}

 

$clang –fopenmp-backend -fopenmp-targets=x86_64-mic, i386-pc-linux-gnu …

 

x86_64-mic

 

define void @__omp_offloading…(i64 %d) #0 {

entry:

}

 

i386-pc-linux-gnu

 

define void @__omp_offloading…(double* dereferenceable(8) %d) #0 {

entry:

 …

}

 

There is an inconsistency between host and target part(s) of the program!

 

Change Made

==========

We proposed new module level attribute to represent target device information.

 

/// Get the target device information, which is a comma-separated string

/// describing one or more devices.

const std::string &getTargetDevices() const { return TargetDevices; }

 

/// set the target device information.

void setTargetDevices(StringRef T) { TargetDevices = T; }

 

IR Dump (the extension indicated in red font)

target triple = "x86_64-unknown-linux-gnu"

target device_triples = "x86_64-mic,i386-pc-linux-gnu"

 

========

Summary:

========

We propose a new module-level attribute to represent the target device information embedded in the LLVM IR. The change is simple and straightforward. The patch is uploaded with this RFC for code review.

https://reviews.llvm.org/D46071

https://reviews.llvm.org/D46074

 

 


_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev
On 4/25/2018 3:05 PM, Lin, Jin via llvm-dev wrote:

RFC: Representing the target device information in the LLVM IR

===========================================================================

 

Why this RFC change?

=================

The target device information needs to be passed to the LLVM backend when OpenMP backend outlining is enabled. For example, for multiple target devices, the target compilation has to generate a single host to support all the targets. In order to make sure all the target outlined functions have the same interface, the information of all the target architectures is needed during host and target compilation. In the following example, the firstprivate variable ‘d’ is passed by value under x86_64-mic and passed by reference under i386-pc-linux-gnu. In order to avoid this inconsistency, the compiler needs all the target architecture information so it can find a common interface. In this example, it will change the x86_64-mic interface for ‘d’ to pass by reference.

 

Existing code: 64-bit firstprivate variable

 

void foo() {

 double d = 1.0;

 #pragma omp target firstprivate(d)

 {}

}

 

$clang –fopenmp-backend -fopenmp-targets=x86_64-mic, i386-pc-linux-gnu …

 

x86_64-mic

 

define void @__omp_offloading…(i64 %d) #0 {

entry:

}

 

i386-pc-linux-gnu

 

define void @__omp_offloading…(double* dereferenceable(8) %d) #0 {

entry:

 …

}

 

There is an inconsistency between host and target part(s) of the program!


I don't see how this inconsistency is a problem... at least, not on its own.  The host code doesn't call either of these functions directly; it calls the OpenMP runtime, which should invoke the offloaded function correctly.  (If it doesn't, that's a bug in the OpenMP lowering, not the LLVM backend.)

-Eli
-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8

  …

  %20 = call i32 @__tgt_target(i64 -1, i8* @.omp_offload.region_id, i32 4, i8** %4, i8** %6, i64* getelementptr inbounds ([4 x i64], [4 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([4 x i64], [4 x i64]* @.offload_maptypes, i32 0, i32 0)) 

 

Thanks,

 

Jin

 

From: Friedman, Eli [mailto:[hidden email]]
Sent: Wednesday, April 25, 2018 3:18 PM
To: Lin, Jin <[hidden email]>; '[hidden email]' <[hidden email]>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:05 PM, Lin, Jin via llvm-dev wrote:

RFC: Representing the target device information in the LLVM IR

===========================================================================

 

Why this RFC change?

=================

The target device information needs to be passed to the LLVM backend when OpenMP backend outlining is enabled. For example, for multiple target devices, the target compilation has to generate a single host to support all the targets. In order to make sure all the target outlined functions have the same interface, the information of all the target architectures is needed during host and target compilation. In the following example, the firstprivate variable ‘d’ is passed by value under x86_64-mic and passed by reference under i386-pc-linux-gnu. In order to avoid this inconsistency, the compiler needs all the target architecture information so it can find a common interface. In this example, it will change the x86_64-mic interface for ‘d’ to pass by reference.

 

Existing code: 64-bit firstprivate variable

 

void foo() {

 double d = 1.0;

 #pragma omp target firstprivate(d)

 {}

}

 

$clang –fopenmp-backend -fopenmp-targets=x86_64-mic, i386-pc-linux-gnu …

 

x86_64-mic

 

define void @__omp_offloading…(i64 %d) #0 {

entry:

}

 

i386-pc-linux-gnu

 

define void @__omp_offloading…(double* dereferenceable(8) %d) #0 {

entry:

 …

}

 

There is an inconsistency between host and target part(s) of the program!


I don't see how this inconsistency is a problem... at least, not on its own.  The host code doesn't call either of these functions directly; it calls the OpenMP runtime, which should invoke the offloaded function correctly.  (If it doesn't, that's a bug in the OpenMP lowering, not the LLVM backend.)

-Eli

-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev
On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli

-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [mailto:[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin <[hidden email]>; '[hidden email]' <[hidden email]>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli


-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev

Hi, Jin,

Can you please back up a bit and talk about the programming environment in which this problem manifests?

If I have a host and a target with different ABIs, then it seems we have lots of problems. For one thing, the layouts of structures are different, the sizes of some integer types are different, the sizes of pointers are different, and so on. It seems like a solution in this space should address, somehow, this general translation problem. Fixing this particular problem with the dispatch function's parameters feels like only the tip of the iceberg. What if I'm passing a pointer to some structure, or a pointer to other pointers, etc.?

I understand that OpenMP v5 is expected to have some custom "mappers" to handle deep copying and translation. Is this related to the design space here?

Thanks again,

Hal


On 04/25/2018 07:22 PM, Lin, Jin via llvm-dev wrote:

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin [hidden email]; '[hidden email]' [hidden email]
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli


-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project


_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev
To follow on here:

This is most assuredly just the tip of the iceberg if you need to co-mingle two targets as part of the module. Basically any solution is going to be better than trying to do that.

-eric

On Wed, Apr 25, 2018 at 5:49 PM Hal Finkel via llvm-dev <[hidden email]> wrote:

Hi, Jin,

Can you please back up a bit and talk about the programming environment in which this problem manifests?

If I have a host and a target with different ABIs, then it seems we have lots of problems. For one thing, the layouts of structures are different, the sizes of some integer types are different, the sizes of pointers are different, and so on. It seems like a solution in this space should address, somehow, this general translation problem. Fixing this particular problem with the dispatch function's parameters feels like only the tip of the iceberg. What if I'm passing a pointer to some structure, or a pointer to other pointers, etc.?

I understand that OpenMP v5 is expected to have some custom "mappers" to handle deep copying and translation. Is this related to the design space here?

Thanks again,

Hal


On 04/25/2018 07:22 PM, Lin, Jin via llvm-dev wrote:

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin [hidden email]; '[hidden email]' [hidden email]
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli


-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project


_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory
_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev
In reply to this post by Dean Michael Berris via llvm-dev

Hi Hal,

   We are not trying to address issues where the object mapped are of different sizes between host and target with different ABI.  The issue is when the objects are of same size like double which is 8bytes on both 32bit and 64bit platform.  If a double is used in a first_private on a target clause,  the 64 bit side will pass it as value whereas on the 32bit side since the value does not fit in the argument it will be passed as pointer to a double. There will be a mis-match at the call site and entry site on this value.

  The main reason for this change is that when we do backend outlining for target pragmas the targets information needs to be communicated to the backend to generate the tables with the right names.  Generate LLVM IR for passing this information is one mechanism and other is passing the command option to the backend.  For the later each pass which needs this info will have to change.

Thanks

Ravi

 

From: Hal Finkel [[hidden email]]
Sent: Wednesday, April 25, 2018 5:50 PM
To: Lin, Jin <[hidden email]>; Friedman, Eli <[hidden email]>; '[hidden email]' <[hidden email]>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

Hi, Jin,

Can you please back up a bit and talk about the programming environment in which this problem manifests?

If I have a host and a target with different ABIs, then it seems we have lots of problems. For one thing, the layouts of structures are different, the sizes of some integer types are different, the sizes of pointers are different, and so on. It seems like a solution in this space should address, somehow, this general translation problem. Fixing this particular problem with the dispatch function's parameters feels like only the tip of the iceberg. What if I'm passing a pointer to some structure, or a pointer to other pointers, etc.?

I understand that OpenMP v5 is expected to have some custom "mappers" to handle deep copying and translation. Is this related to the design space here?

Thanks again,

Hal

 

On 04/25/2018 07:22 PM, Lin, Jin via llvm-dev wrote:

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin [hidden email]; '[hidden email]' [hidden email]
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli


-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project



_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

 

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev


On 04/26/2018 07:03 PM, Narayanaswamy, Ravi wrote:

Hi Hal,

   We are not trying to address issues where the object mapped are of different sizes between host and target with different ABI.


Why are you not trying to address that issue?

 -Hal

The issue is when the objects are of same size like double which is 8bytes on both 32bit and 64bit platform.  If a double is used in a first_private on a target clause,  the 64 bit side will pass it as value whereas on the 32bit side since the value does not fit in the argument it will be passed as pointer to a double. There will be a mis-match at the call site and entry site on this value.

  The main reason for this change is that when we do backend outlining for target pragmas the targets information needs to be communicated to the backend to generate the tables with the right names.  Generate LLVM IR for passing this information is one mechanism and other is passing the command option to the backend.  For the later each pass which needs this info will have to change.

Thanks

Ravi

 

From: Hal Finkel [[hidden email]]
Sent: Wednesday, April 25, 2018 5:50 PM
To: Lin, Jin <[hidden email]>; Friedman, Eli <[hidden email]>; '[hidden email]' <[hidden email]>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

Hi, Jin,

Can you please back up a bit and talk about the programming environment in which this problem manifests?

If I have a host and a target with different ABIs, then it seems we have lots of problems. For one thing, the layouts of structures are different, the sizes of some integer types are different, the sizes of pointers are different, and so on. It seems like a solution in this space should address, somehow, this general translation problem. Fixing this particular problem with the dispatch function's parameters feels like only the tip of the iceberg. What if I'm passing a pointer to some structure, or a pointer to other pointers, etc.?

I understand that OpenMP v5 is expected to have some custom "mappers" to handle deep copying and translation. Is this related to the design space here?

Thanks again,

Hal

 

On 04/25/2018 07:22 PM, Lin, Jin via llvm-dev wrote:

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin [hidden email]; '[hidden email]' [hidden email]
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli


-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project



_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

 

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev

I think because it is prohibited by the standard. According to OpenMP standard the new copies of the variables must be of the same type and the same size. If the type has the different size on the device, we become not compatible with the standard.

I agree with Eric Christofer. From my point of view, it just breaks the existing ABI. Instead of breaking the existing ABI, I think, it would better to introduce a new, portable, ABI. The classic ABI could be used to get a little bit more performance, while the new one could be used to get the compatibility between pointer-size incompatible targets. In this ABI all variables must be passed by reference. I think it will solve all the problems.

With the proposed change there would be a difference between the code compiled for (1) 64 bit host+64 bit device and (2) 64 bit host +64|32 bit device. You won't be able to make it work properly the program linked from the (2)nd host object + the (1)st device object because of the different ABIs. I think it is better to explicitly specify that we're going to use some special ABI rather than doing such tricky and dangerous things as in this proposal.
-------------
Best regards,
Alexey Bataev
26.04.2018 20:14, Hal Finkel via llvm-dev пишет:


On 04/26/2018 07:03 PM, Narayanaswamy, Ravi wrote:

Hi Hal,

   We are not trying to address issues where the object mapped are of different sizes between host and target with different ABI.


Why are you not trying to address that issue?

 -Hal

The issue is when the objects are of same size like double which is 8bytes on both 32bit and 64bit platform.  If a double is used in a first_private on a target clause,  the 64 bit side will pass it as value whereas on the 32bit side since the value does not fit in the argument it will be passed as pointer to a double. There will be a mis-match at the call site and entry site on this value.

  The main reason for this change is that when we do backend outlining for target pragmas the targets information needs to be communicated to the backend to generate the tables with the right names.  Generate LLVM IR for passing this information is one mechanism and other is passing the command option to the backend.  For the later each pass which needs this info will have to change.

Thanks

Ravi

 

From: Hal Finkel [[hidden email]]
Sent: Wednesday, April 25, 2018 5:50 PM
To: Lin, Jin <[hidden email]>; Friedman, Eli <[hidden email]>; '[hidden email]' <[hidden email]>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

Hi, Jin,

Can you please back up a bit and talk about the programming environment in which this problem manifests?

If I have a host and a target with different ABIs, then it seems we have lots of problems. For one thing, the layouts of structures are different, the sizes of some integer types are different, the sizes of pointers are different, and so on. It seems like a solution in this space should address, somehow, this general translation problem. Fixing this particular problem with the dispatch function's parameters feels like only the tip of the iceberg. What if I'm passing a pointer to some structure, or a pointer to other pointers, etc.?

I understand that OpenMP v5 is expected to have some custom "mappers" to handle deep copying and translation. Is this related to the design space here?

Thanks again,

Hal

 

On 04/25/2018 07:22 PM, Lin, Jin via llvm-dev wrote:

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin [hidden email]; '[hidden email]' [hidden email]
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli


-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project



_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

 

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory


_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

signature.asc (849 bytes) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

Dean Michael Berris via llvm-dev
In reply to this post by Dean Michael Berris via llvm-dev

 

On 04/26/2018 07:03 PM, Narayanaswamy, Ravi wrote:

Hi Hal,

   We are not trying to address issues where the object mapped are of different sizes between host and target with different ABI.


Why are you not trying to address that issue?

 

Good question.  It is a harder problem and no vendor except Intel (offloading from windows to Linux) required this and so were not pushing for this.

 

Also bring up a whole lot of problems with shared memory, how do use shared memory if the objects are different, are copies made, when is the original item synced up.

 

Ravi




The issue is when the objects are of same size like double which is 8bytes on both 32bit and 64bit platform.  If a double is used in a first_private on a target clause,  the 64 bit side will pass it as value whereas on the 32bit side since the value does not fit in the argument it will be passed as pointer to a double. There will be a mis-match at the call site and entry site on this value.

  The main reason for this change is that when we do backend outlining for target pragmas the targets information needs to be communicated to the backend to generate the tables with the right names.  Generate LLVM IR for passing this information is one mechanism and other is passing the command option to the backend.  For the later each pass which needs this info will have to change.

Thanks

Ravi

 

From: Hal Finkel [[hidden email]]
Sent: Wednesday, April 25, 2018 5:50 PM
To: Lin, Jin <[hidden email]>; Friedman, Eli <[hidden email]>; '[hidden email]' <[hidden email]>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

Hi, Jin,

Can you please back up a bit and talk about the programming environment in which this problem manifests?

If I have a host and a target with different ABIs, then it seems we have lots of problems. For one thing, the layouts of structures are different, the sizes of some integer types are different, the sizes of pointers are different, and so on. It seems like a solution in this space should address, somehow, this general translation problem. Fixing this particular problem with the dispatch function's parameters feels like only the tip of the iceberg. What if I'm passing a pointer to some structure, or a pointer to other pointers, etc.?

I understand that OpenMP v5 is expected to have some custom "mappers" to handle deep copying and translation. Is this related to the design space here?

Thanks again,

Hal

 

On 04/25/2018 07:22 PM, Lin, Jin via llvm-dev wrote:

For the firstprivate clause, the compiler generates code to pass it  by value or by reference to the outlined function. The reason the first private scalars is generally passed by value is for the performance reason.

For this particular case, the compiler cannot generate code to pass the double @gg by value under i386-pc-linux-gnu since the value is 64 bit while the architecture is 32bit.

For the host compilation, the compiler generates the code to pass the data as well as the outlined function name to the OMP runtime.

For the target compilation, the compiler generates the outlined function so that it can be called by the OMP runtime. 

So, the compiler is required to generate a single call on the host to support all the targets. All the target versions must have the same interface. So the common interface of the outline function should be used. For this particular example, the variable @gcc should be passed by reference under x86_64-mic.

Please let me know if you have more questions.

Jin

 

From: Friedman, Eli [[hidden email]]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin [hidden email]; '[hidden email]' [hidden email]
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

 

On 4/25/2018 3:48 PM, Lin, Jin wrote:

Given a global variable @gg, the compiler has to generate code on the host to specify whether it is passed by value or passed by reference. In the following example, if the compiler generates the code for passing by value, the outlined function on the target i386-pc-linux-gnu cannot get the correct value since it assumes the variable @gg is passed by reference.  

 

Here is the corresponding IR on the host side.

  %0 = load double, double* @gg, align 8, !tbaa !3

  %1 = bitcast double %0 to i64

   …

  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0, i32 2

  %13 = bitcast i8** %12 to i64*

  store i64 %1, i64* %13, align 8


Could you describe the overall process of calling an offloaded function in a bit more detail?  How do you describe the ABI of the called function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers into offload_baseptrs.

-Eli



-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project




_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev

 

-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory



-- 
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory

_______________________________________________
LLVM Developers mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev