GSoC Interested student

Greetings,

I'm an LLVM newcomer interested in participating in the next GSoC under the project "Improve parallelism-aware analyses and optimizations". I've already read "Compiler Optimizations for OpenMP" and "Compiler Optimizations for Parallel Programs" both by Doerfert and Finkel. Also, I've watched their two LLVM meeting developers conferences "Representing parallelism within LLVM" and "Optimizing indirections, using abstractions without ...", so I have some background in the problems they are trying to tackle. Also, I'm close to solve one of the bugs posted in bugzilla, which has given me a broader idea of how the OpenMP pragmas are translated. The bus is this: 40253 – combined target teams implements shared clause as firstprivate.

I tried contacting you Johannes but with your Argonne email, I suppose that's why you couldn't answer. Could you tell me more about the project?, what variants do you have?, or what can I focus my application on?. Should I start with what you recommended to Emanuel?.

Thank you.

Hi Hamilton,

Thanks for your interest in LLVM!

IMHO you have a good start. I’ll try to help by CC’ing Johannes. Note that it’s difficult for LLVM developers
(including GSoC mentors and including Johannes) to keep up with every discussion on llvm-dev.
So, you can CC them to make their life easier. :slight_smile:

Johannes, I see a lot of people interested in this project. Maybe a list of starting points, patches etc.
could help.

Best,
Stefanos

Στις Τρί, 10 Μαρ 2020 στις 1:04 π.μ., ο/η Hamilton Tobon Mosquera via llvm-dev <llvm-dev@lists.llvm.org> έγραψε:

Hi there!,

As part of my application process to the next GSoC I’m working on the TODO in OpenMPOpt.cpp line 437:

// TODO: We should validate the declaration agains the types we expect.

Link:

I have a question. When there is a mismatch in the types (return type or argument types) between the function declaration found and the runtime library function, what should I do?. Show an error or a warning?, continuing with the pass or exiting the program?.

Thanks.

Hi Hamilton,

first, apologies for not responding to your initial mail. (I got it on
the anl.gov email but just forgot the reply, ...)

I'm glad you took the initiative to work on a bug and TODO!

Wrt. your question, I guess we should just ignore the runtime function
pretending it is something different (for now). It should not happen to
often (I hope) and this is more of a protection agains accidental
miscompilation when people experiment with the runtime.

@Stefanos
Correct, we need more structure. I started a dashboard but I got
distracted. I will create a google doc with tasks today and share it
with people.

Cheers,
  Johannes

Hi there!,

I think I’m done with this patch: D76058. If that is everything concerning that, could you please assign me another task.

I’ve been reading about Concurrent Control Flow Graphs, which would be awesome for LLVM because it would be kind of a starting point for optimizations for parallel programs in general. The thing is that it seems more complicated than it is presented in the papers, I think that’s why it’s not implemented in LLVM yet. So, I think I’ll try searching for something else. Nevertheless, don’t you think that implementing this is worth the effort?. The papers I’m reading are from 1990s, I don’t understand why they aren’t widely implemented in big compilers xd.

I saw this commit from you: . Is there something I can help with concerning this?. I would like to do contribute to the OpenMP target offloading. Is there anything I can help with?.

Hi Hamilton,

I personally don't believe the effort to make the IR "parallelism-aware" is worth it right now.
What I propose is something else, see for example 2018 LLVM Developers’ Meeting: J. Doerfert “ Optimizing Indirections, using abstractions without...” - YouTube and the OpenMPOpt pass.
There are multiple bigger tasks towards better offloading, one of which is described here [OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub
It might be interesting to add the transfer functions used for memory transfer to the OMPKinds.def file so they are known in the OpenMPOpt.cpp pass.
Another small task would be to eliminate trivially redundant ones, that is if you have copy X from A to B directly followed by copy X from B to A, remove both.

Does this make sense?

Cheers,
  Johannes

Hi there!,

Right now I'm starting with the project proposed by Johannes here: [OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub.

The first task is identifying OpenMP API calls referred to device memory allocation. I first understood how the compiler driver orchestrates the process of offloading OpenMP regions to devices, and now I'm looking at the step of the pipeline that generates LLVM IR for both host and device. I wrote a simple C program that uses OpenMP to map a local array to the device main memory. Then I'm looking at the device and host LLVM IR representation of this program, but I can't find any runtime call referred to device memory allocation. In the host IR I see that the array is passed to the outlined function and stored in a register inside the outlined function, but there's no memory allocation runtime call. And in the device IR the array is passed to the main offloading function. Any suggestion or documentation on how that memory transfer is achieved?, and how to track its API calls. It might be that I'm not using the proper omp directive, so here is C code:

#include <omp.h>
#include <stdio.h>

#define size 1000000
double a[size];

int main() {

#pragma omp target map(to: a[0:size])
#pragma omp teams distribute
for (int i = 0; i < size; ++i) {
a[i] = a[i]*(i*i)/2;
}
return 0;

PD: The command to compile is this: clang -v -save-temps -I/path/to/llvm/release/9.x/include -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_61 memory_transfer.c -o memory_transfer

Thank you all.

Hamilton.

Hi there Johannes!,

I was checking out the issue (about latency hiding when transferring memory from host to device) you posted about a week ago, but it's not there anymore. Did you remove it?.

Thanks.

Hamilton.

I did not. The issues sections on GH was removed.

Ahhh I understand. So I believe the project you posted is still opened. Do you any advice on the question I posted before?.

Thank you.

Ahhh I understand. So I believe the project you posted is still

opened. Do you any advice on the question I posted before?.

Yes. Still open. You should use this as the main part of the GSoC
proposal (if you are interested). Can you repeat the question?

>
> Thank you.
>
>> I did not. The issues sections on GH was removed.
>>
>> ________________________________________
>> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> Sent: Sunday, March 22, 2020 12:21
>> To: llvm-dev@lists.llvm.org
>> Cc: Doerfert, Johannes
>> Subject: Re: [llvm-dev] GSoC Interested student
>>
>> Hi there Johannes!,
>>
>> I was checking out the issue (about latency hiding when transferring
>> memory from host to device) you posted about a week ago, but it's not
>> there anymore. Did you remove it?.
>>
>> Thanks.
>>
>> Hamilton.
>>
>>> Hi there!,
>>>
>>> Right now I'm starting with the project proposed by Johannes here:
>>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub.
>>>
>>>
>>> The first task is identifying OpenMP API calls referred to device
memory
>>> allocation. I first understood how the compiler driver orchestrates the
>>> process of offloading OpenMP regions to devices, and now I'm looking at
>>> the step of the pipeline that generates LLVM IR for both host and
>>> device. I wrote a simple C program that uses OpenMP to map a local
array
>>> to the device main memory. Then I'm looking at the device and host LLVM
>>> IR representation of this program, but I can't find any runtime call
>>> referred to device memory allocation. In the host IR I see that the
>>> array is passed to the outlined function and stored in a register
inside
>>> the outlined function, but there's no memory allocation runtime call.
>>> And in the device IR the array is passed to the main offloading
>>> function. Any suggestion or documentation on how that memory
transfer is
>>> achieved?, and how to track its API calls. It might be that I'm not
>>> using the proper omp directive, so here is C code:
>>>
>>>
>>> #include <omp.h>
>>> #include <stdio.h>
>>>
>>> #define size 1000000
>>> double a[size];
>>>
>>> int main() {
>>>
>>> #pragma omp target map(to: a[0:size])
>>> #pragma omp teams distribute
>>> for (int i = 0; i < size; ++i) {
>>> a[i] = a[i]*(i*i)/2;
>>> }
>>> return 0;
>>>
>>> PD: The command to compile is this: clang -v -save-temps
>>> -I/path/to/llvm/release/9.x/include -fopenmp
>>> -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_61
>>> memory_transfer.c -o memory_transfer
>>>
>>> Thank you all.
>>>
>>> Hamilton.
>>>
>>>
>>>> Hi Hamilton,
>>>>
>>>> I personally don't believe the effort to make the IR
>>>> "parallelism-aware" is worth it right now.
>>>> What I propose is something else, see for example
>>>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fyoutu.be%2FzfiHaPaoQPc&amp;data=01|01|htobonm%40eafit.edu.co|ec2747b9c96d4b8ddc0f08d7ce862a84|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=fXYu60Fv7uhvW%2BY6%2F7Rw%2FAfXQEXdDbRkiRaf4KfVNgs%3D&amp;reserved=0
>>>> and the OpenMPOpt pass.
>>>> There are multiple bigger tasks towards better offloading, one of
>>>> which is described here
>>>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub
>>>> It might be interesting to add the transfer functions used for memory
>>>> transfer to the OMPKinds.def file so they are known in the
>>>> OpenMPOpt.cpp pass.
>>>> Another small task would be to eliminate trivially redundant ones,
>>>> that is if you have copy X from A to B directly followed by copy X
>>>> from B to A, remove both.
>>>>
>>>> Does this make sense?
>>>>
>>>> Cheers,
>>>> Johannes
>>>>
>>>>
>>>>
>>>>
>>>>
>>>> ---------------------------------------
>>>> Johannes Doerfert
>>>> Researcher
>>>>
>>>> Argonne National Laboratory
>>>> Lemont, IL 60439, USA
>>>>
>>>> jdoerfert@anl.gov
>>>>
>>>> ________________________________________
>>>> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>>>> Sent: Wednesday, March 18, 2020 19:14
>>>> To: llvm-dev@lists.llvm.org
>>>> Cc: Doerfert, Johannes
>>>> Subject: Re: [llvm-dev] GSoC Interested student
>>>>
>>>> Hi there!,
>>>>
>>>> I think I'm done with this patch: D76058. If that is everything
>>>> concerning that, could you please assign me another task.
>>>>
>>>> I've been reading about Concurrent Control Flow Graphs, which would
>>>> be awesome for LLVM because it would be kind of a starting point for
>>>> optimizations for parallel programs in general. The thing is that it
>>>> seems more complicated than it is presented in the papers, I think
>>>> that's why it's not implemented in LLVM yet. So, I think I'll try
>>>> searching for something else. Nevertheless, don't you think that
>>>> implementing this is worth the effort?. The papers I'm reading are
>>>> from 1990s, I don't understand why they aren't widely implemented in
>>>> big compilers xd.
>>>>
>>>> I saw this commit from you:
>>>>
⚙ D29250 [PIR] Add fork, join and halt instructions.
>>>> Is there something I can help with concerning this?.
>>>>
>>>> I would like to do contribute to the OpenMP target offloading. Is
>>>> there anything I can help with?.
>>>>
>>>>
>>>> Hi there!,
>>>>
>>>> As part of my application process to the next GSoC I'm working on the
>>>> TODO in OpenMPOpt.cpp line 437:
>>>>
>>>> // TODO: We should validate the declaration agains the types we
>>>> expect.
>>>>
>>>> Link:
>>>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fblob%2Fmaster%2Fllvm%2Flib%2FTransforms%2FIPO%2FOpenMPOpt.cpp%23L437&amp;data=01|01|htobonm%40eafit.edu.co|ec2747b9c96d4b8ddc0f08d7ce862a84|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=fo5aGg50Jj9qDj9WmJuwCUbwqxUTVJZsBsWb2UBPSVw%3D&amp;reserved=0
>>>>
>>>> I have a question. When there is a mismatch in the types (return type
>>>> or argument types) between the function declaration found and the
>>>> runtime library function, what should I do?. Show an error or a
>>>> warning?, continuing with the pass or exiting the program?.
>>>>
>>>> Thanks.
>>>>
>>>>
>>>> Hi Hamilton,
>>>>
>>>> Thanks for your interest in LLVM!
>>>>
>>>> IMHO you have a good start. I'll try to help by CC'ing Johannes. Note
>>>> that it's difficult for LLVM developers
>>>> (including GSoC mentors and including Johannes) to keep up with every
>>>> discussion on llvm-dev.
>>>> So, you can CC them to make their life easier. :slight_smile:
>>>>
>>>> Johannes, I see a lot of people interested in this project. Maybe a
>>>> list of starting points, patches etc.
>>>> could help.
>>>>
>>>> Best,
>>>> Stefanos
>>>>
>>>> Στις Τρί, 10 Μαρ 2020 στις 1:04 π.μ., ο/η Hamilton Tobon Mosquera via
>>>> llvm-dev <llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>>
>>>> έγραψε:
>>>> Greetings,
>>>>
>>>>
>>>> I'm an LLVM newcomer interested in participating in the next GSoC
under
>>>> the project "Improve parallelism-aware analyses and
optimizations". I've
>>>> already read "Compiler Optimizations for OpenMP" and "Compiler
>>>> Optimizations for Parallel Programs" both by Doerfert and Finkel.
Also,
>>>> I've watched their two LLVM meeting developers conferences
"Representing
>>>> parallelism within LLVM" and "Optimizing indirections, using
>>>> abstractions without ...", so I have some background in the problems
>>>> they are trying to tackle. Also, I'm close to solve one of the bugs
>>>> posted in bugzilla, which has given me a broader idea of how the
OpenMP
>>>> pragmas are translated. The bus is this:
>>>>
40253 – combined target teams implements shared clause as firstprivate.
>>>>
>>>>
>>>> I tried contacting you Johannes but with your Argonne email, I suppose
>>>> that's why you couldn't answer. Could you tell me more about the
>>>> project?, what variants do you have?, or what can I focus my
application
>>>> on?. Should I start with what you recommended to Emanuel?.
>>>>
>>>> Thank you.
>>>>
>>>> _______________________________________________
>>>> LLVM Developers mailing list
>>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>>>>
llvm-dev Info Page
>>>>
>>>> La información contenida en este correo electrónico está dirigida
>>>> únicamente a su destinatario y puede contener información
>>>> confidencial, material privilegiado o información protegida por
>>>> derecho de autor. Está prohibida cualquier copia, utilización,
>>>> indebida retención, modificación, difusión, distribución o
>>>> reproducción total o parcial. Si usted recibe este mensaje por error,
>>>> por favor contacte al remitente y elimínelo. La información aquí
>>>> contenida es responsabilidad exclusiva de su remitente por lo tanto
>>>> la Universidad EAFIT no se hace responsable de lo que el mensaje
>>>> contenga. The information contained in this email is addressed to its
>>>> recipient only and may contain confidential information, privileged
>>>> material or information protected by copyright. Its prohibited any
>>>> copy, use, improper retention, modification, dissemination,
>>>> distribution or total or partial reproduction. If you receive this
>>>> message by error, please contact the sender and delete it. The
>>>> information contained herein is the sole responsibility of the sender
>>>> therefore Universidad EAFIT is not responsible for what the message
>>>> contains.
>>>> La información contenida en este correo electrónico está dirigida
>>>> únicamente a su destinatario y puede contener información
>>>> confidencial, material privilegiado o información protegida por
>>>> derecho de autor. Está prohibida cualquier copia, utilización,
>>>> indebida retención, modificación, difusión, distribución o
>>>> reproducción total o parcial. Si usted recibe este mensaje por error,
>>>> por favor contacte al remitente y elimínelo. La información aquí
>>>> contenida es responsabilidad exclusiva de su remitente por lo tanto
>>>> la Universidad EAFIT no se hace responsable de lo que el mensaje
>>>> contenga. The information contained in this email is addressed to its
>>>> recipient only and may contain confidential information, privileged
>>>> material or information protected by copyright. Its prohibited any
>>>> copy, use, improper retention, modification, dissemination,
>>>> distribution or total or partial reproduction. If you receive this
>>>> message by error, please contact the sender and delete it. The
>>>> information contained herein is the sole responsibility of the sender
>>>> therefore Universidad EAFIT is not responsible for what the message
>>>> contains.
>>>>
>>>>
>>>> _______________________________________________
>>>> LLVM Developers mailing list
>>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>>>>
llvm-dev Info Page
>>>>
>>>>
>>>> La información contenida en este correo electrónico está dirigida
>>>> únicamente a su destinatario y puede contener información
>>>> confidencial, material privilegiado o información protegida por
>>>> derecho de autor. Está prohibida cualquier copia, utilización,
>>>> indebida retención, modificación, difusión, distribución o
>>>> reproducción total o parcial. Si usted recibe este mensaje por error,
>>>> por favor contacte al remitente y elimínelo. La información aquí
>>>> contenida es responsabilidad exclusiva de su remitente por lo tanto
>>>> la Universidad EAFIT no se hace responsable de lo que el mensaje
>>>> contenga. The information contained in this email is addressed to its
>>>> recipient only and may contain confidential information, privileged
>>>> material or information protected by copyright. Its prohibited any
>>>> copy, use, improper retention, modification, dissemination,
>>>> distribution or total or partial reproduction. If you receive this
>>>> message by error, please contact the sender and delete it. The
>>>> information contained herein is the sole responsibility of the sender
>>>> therefore Universidad EAFIT is not responsible for what the message
>>>> contains.
>>> _______________________________________________
>>> LLVM Developers mailing list
>>> llvm-dev@lists.llvm.org
>>>
llvm-dev Info Page
>>>
>>> La información contenida en este correo electrónico está dirigida
>>> únicamente a su destinatario y puede contener información
>>> confidencial, material privilegiado o información protegida por
>>> derecho de autor. Está prohibida cualquier copia, utilización,
>>> indebida retención, modificación, difusión, distribución o
>>> reproducción total o parcial. Si usted recibe este mensaje por error,
>>> por favor contacte al remitente y elimínelo. La información aquí
>>> contenida es responsabilidad exclusiva de su remitente por lo tanto la
>>> Universidad EAFIT no se hace responsable de lo que el mensaje
>>> contenga. The information contained in this email is addressed to its
>>> recipient only and may contain confidential information, privileged
>>> material or information protected by copyright. Its prohibited any
>>> copy, use, improper retention, modification, dissemination,
>>> distribution or total or partial reproduction. If you receive this
>>> message by error, please contact the sender and delete it. The
>>> information contained herein is the sole responsibility of the sender
>>> therefore Universidad EAFIT is not responsible for what the message
>>> contains.
>> La información contenida en este correo electrónico está dirigida
únicamente a su destinatario y puede contener información confidencial,
material privilegiado o información protegida por derecho de autor. Está
prohibida cualquier copia, utilización, indebida retención,
modificación, difusión, distribución o reproducción total o parcial. Si
usted recibe este mensaje por error, por favor contacte al remitente y
elimínelo. La información aquí contenida es responsabilidad exclusiva de
su remitente por lo tanto la Universidad EAFIT no se hace responsable de
lo que el mensaje contenga. The information contained in this email is
addressed to its recipient only and may contain confidential
information, privileged material or information protected by copyright.
Its prohibited any copy, use, improper retention, modification,
dissemination, distribution or total or partial reproduction. If you
receive this message by error, please contact the sender and delete it.
The information contained herein is the sole responsibility of the
sender therefore Universidad EAFIT is not responsible for what the
message contains.

Hi Johannes,
Yes, I'm highly interested, I already started. It would be nice if you post the details again. I didn't take note of them, I was always looking at the issue in GH.

The problem I'm facing right now is this:
The first task is identifying OpenMP API calls referred to device memory
allocation. I first understood how the compiler driver orchestrates the
process of offloading OpenMP regions to devices, and now I'm looking at
the step of the pipeline that generates LLVM IR for both host and
device. I wrote a simple C program that uses OpenMP to map a local array
to the device main memory. Then I'm looking at the device and host LLVM
IR representation of this program, but I can't find any runtime call
referred to device memory allocation. In the host IR I see that the
array is passed to the outlined function and stored in a register inside
the outlined function, but there's no memory allocation runtime call.
And in the device IR the array is passed to the main offloading
function.
Any suggestion or documentation on how that memory transfer is
achieved?, and how to track its API calls. It might be that I'm not
using the proper omp directive, so here is C code:

#include <omp.h>
#include <stdio.h>

#define size 1000000
double a[size];

int main() {

#pragma omp target map(to: a[0:size])
#pragma omp teams distribute
for (int i = 0; i < size; ++i) {
a[i] = a[i]*(i*i)/2;
}

return 0;
}

PD: The command to compile is this:
  clang -v -save-temps \
  -I/path/to/llvm/release/9.x/include -fopenmp \
  -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
         -march=sm_61 memory_transfer.c -o memory_transfer

Thanks.
Hamilton.

Hi Johannes,

>
> Yes, I'm highly interested, I already started. It would be nice if you
> post the details again. I didn't take note of them, I was always
> looking at the issue in GH.

I asked if we can make the readonly but unsure yet. I can write it again
but it'll take ma some time.

> The problem I'm facing right now is this:
> The first task is identifying OpenMP API calls referred to device memory
> allocation. I first understood how the compiler driver orchestrates the
> process of offloading OpenMP regions to devices, and now I'm looking at
> the step of the pipeline that generates LLVM IR for both host and
> device. I wrote a simple C program that uses OpenMP to map a local array
> to the device main memory. Then I'm looking at the device and host LLVM
> IR representation of this program, but I can't find any runtime call
> referred to device memory allocation. In the host IR I see that the
> array is passed to the outlined function and stored in a register inside
> the outlined function, but there's no memory allocation runtime call.
> And in the device IR the array is passed to the main offloading
> function.
> Any suggestion or documentation on how that memory transfer is
> achieved?, and how to track its API calls. It might be that I'm not
> using the proper omp directive, so here is C code:

Take a look at Compiler Explorer , the mapping request,

which may or may not cause data-transfers, are part of the call to
__tgt_target_teams[_nowait] and __tgt_target_data_update[_nowait]. Start
by reading the impl. of those two functions in the openmp runtime
(llvm-project/openmp/...).

Let me know if this helps and once you have new questions.

Cheers,
Johannes

P.S. I merged your first patch. You didn't run all test before. I did
and noticed that we failed one because of type mismatches. Fixed them
before I merged your patch.

>
> #include <omp.h>
> #include <stdio.h>
>
> #define size 1000000
> double a[size];
>
> int main() {
>
> #pragma omp target map(to: a[0:size])
> #pragma omp teams distribute
> for (int i = 0; i < size; ++i) {
> a[i] = a[i]*(i*i)/2;
> }
>
> return 0;
> }
>
>
> PD: The command to compile is this:
> clang -v -save-temps \
> -I/path/to/llvm/release/9.x/include -fopenmp \
> -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
> -march=sm_61 memory_transfer.c -o memory_transfer
>
> Thanks.
> Hamilton.
>
>> > Ahhh I understand. So I believe the project you posted is still
>> opened. Do you any advice on the question I posted before?.
>>
>> Yes. Still open. You should use this as the main part of the GSoC
>> proposal (if you are interested). Can you repeat the question?
>>
>> >
>> > Thank you.
>> >
>> >> I did not. The issues sections on GH was removed.
>> >>
>> >> ________________________________________
>> >> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >> Sent: Sunday, March 22, 2020 12:21
>> >> To: llvm-dev@lists.llvm.org
>> >> Cc: Doerfert, Johannes
>> >> Subject: Re: [llvm-dev] GSoC Interested student
>> >>
>> >> Hi there Johannes!,
>> >>
>> >> I was checking out the issue (about latency hiding when
transferring
>> >> memory from host to device) you posted about a week ago, but
it's not
>> >> there anymore. Did you remove it?.
>> >>
>> >> Thanks.
>> >>
>> >> Hamilton.
>> >>
>> >>> Hi there!,
>> >>>
>> >>> Right now I'm starting with the project proposed by Johannes here:
>> >>>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub.
>> >>>
>> >>>
>> >>> The first task is identifying OpenMP API calls referred to device
>> memory
>> >>> allocation. I first understood how the compiler driver
orchestrates the
>> >>> process of offloading OpenMP regions to devices, and now I'm
looking at
>> >>> the step of the pipeline that generates LLVM IR for both host and
>> >>> device. I wrote a simple C program that uses OpenMP to map a local
>> array
>> >>> to the device main memory. Then I'm looking at the device and
host LLVM
>> >>> IR representation of this program, but I can't find any
runtime call
>> >>> referred to device memory allocation. In the host IR I see
that the
>> >>> array is passed to the outlined function and stored in a register
>> inside
>> >>> the outlined function, but there's no memory allocation
runtime call.
>> >>> And in the device IR the array is passed to the main offloading
>> >>> function. Any suggestion or documentation on how that memory
>> transfer is
>> >>> achieved?, and how to track its API calls. It might be that
I'm not
>> >>> using the proper omp directive, so here is C code:
>> >>>
>> >>>
>> >>> #include <omp.h>
>> >>> #include <stdio.h>
>> >>>
>> >>> #define size 1000000
>> >>> double a[size];
>> >>>
>> >>> int main() {
>> >>>
>> >>> #pragma omp target map(to: a[0:size])
>> >>> #pragma omp teams distribute
>> >>> for (int i = 0; i < size; ++i) {
>> >>> a[i] = a[i]*(i*i)/2;
>> >>> }
>> >>> return 0;
>> >>>
>> >>> PD: The command to compile is this: clang -v -save-temps
>> >>> -I/path/to/llvm/release/9.x/include -fopenmp
>> >>> -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target -march=sm_61
>> >>> memory_transfer.c -o memory_transfer
>> >>>
>> >>> Thank you all.
>> >>>
>> >>> Hamilton.
>> >>>
>> >>>
>> >>>> Hi Hamilton,
>> >>>>
>> >>>> I personally don't believe the effort to make the IR
>> >>>> "parallelism-aware" is worth it right now.
>> >>>> What I propose is something else, see for example
>> >>>>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fyoutu.be%2FzfiHaPaoQPc&amp;data=01|01|htobonm%40eafit.edu.co|27650f5cf4e24c5920be08d7cf3c8209|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=6nejCYt05TQdtstWkksVMf2%2FZyarC3GLEifC0Iyw%2Bic%3D&amp;reserved=0
>> >>>> and the OpenMPOpt pass.
>> >>>> There are multiple bigger tasks towards better offloading, one of
>> >>>> which is described here
>> >>>>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub
>> >>>> It might be interesting to add the transfer functions used
for memory
>> >>>> transfer to the OMPKinds.def file so they are known in the
>> >>>> OpenMPOpt.cpp pass.
>> >>>> Another small task would be to eliminate trivially redundant
ones,
>> >>>> that is if you have copy X from A to B directly followed by
copy X
>> >>>> from B to A, remove both.
>> >>>>
>> >>>> Does this make sense?
>> >>>>
>> >>>> Cheers,
>> >>>> Johannes
>> >>>>
>> >>>>
>> >>>>
>> >>>>
>> >>>>
>> >>>> ---------------------------------------
>> >>>> Johannes Doerfert
>> >>>> Researcher
>> >>>>
>> >>>> Argonne National Laboratory
>> >>>> Lemont, IL 60439, USA
>> >>>>
>> >>>> jdoerfert@anl.gov
>> >>>>
>> >>>> ________________________________________
>> >>>> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >>>> Sent: Wednesday, March 18, 2020 19:14
>> >>>> To: llvm-dev@lists.llvm.org
>> >>>> Cc: Doerfert, Johannes
>> >>>> Subject: Re: [llvm-dev] GSoC Interested student
>> >>>>
>> >>>> Hi there!,
>> >>>>
>> >>>> I think I'm done with this patch: D76058. If that is everything
>> >>>> concerning that, could you please assign me another task.
>> >>>>
>> >>>> I've been reading about Concurrent Control Flow Graphs, which
would
>> >>>> be awesome for LLVM because it would be kind of a starting
point for
>> >>>> optimizations for parallel programs in general. The thing is
that it
>> >>>> seems more complicated than it is presented in the papers, I
think
>> >>>> that's why it's not implemented in LLVM yet. So, I think I'll try
>> >>>> searching for something else. Nevertheless, don't you think that
>> >>>> implementing this is worth the effort?. The papers I'm
reading are
>> >>>> from 1990s, I don't understand why they aren't widely
implemented in
>> >>>> big compilers xd.
>> >>>>
>> >>>> I saw this commit from you:
>> >>>>
>>
⚙ D29250 [PIR] Add fork, join and halt instructions.
>> >>>> Is there something I can help with concerning this?.
>> >>>>
>> >>>> I would like to do contribute to the OpenMP target offloading. Is
>> >>>> there anything I can help with?.
>> >>>>
>> >>>>
>> >>>> Hi there!,
>> >>>>
>> >>>> As part of my application process to the next GSoC I'm
working on the
>> >>>> TODO in OpenMPOpt.cpp line 437:
>> >>>>
>> >>>> // TODO: We should validate the declaration agains the
types we
>> >>>> expect.
>> >>>>
>> >>>> Link:
>> >>>>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fblob%2Fmaster%2Fllvm%2Flib%2FTransforms%2FIPO%2FOpenMPOpt.cpp%23L437&amp;data=01|01|htobonm%40eafit.edu.co|27650f5cf4e24c5920be08d7cf3c8209|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=SbO3ePJdrdDaA1RjUPgOzS7vplRKOe8ZjY68DeM66NQ%3D&amp;reserved=0
>> >>>>
>> >>>> I have a question. When there is a mismatch in the types
(return type
>> >>>> or argument types) between the function declaration found and the
>> >>>> runtime library function, what should I do?. Show an error or a
>> >>>> warning?, continuing with the pass or exiting the program?.
>> >>>>
>> >>>> Thanks.
>> >>>>
>> >>>>
>> >>>> Hi Hamilton,
>> >>>>
>> >>>> Thanks for your interest in LLVM!
>> >>>>
>> >>>> IMHO you have a good start. I'll try to help by CC'ing
Johannes. Note
>> >>>> that it's difficult for LLVM developers
>> >>>> (including GSoC mentors and including Johannes) to keep up
with every
>> >>>> discussion on llvm-dev.
>> >>>> So, you can CC them to make their life easier. :slight_smile:
>> >>>>
>> >>>> Johannes, I see a lot of people interested in this project.
Maybe a
>> >>>> list of starting points, patches etc.
>> >>>> could help.
>> >>>>
>> >>>> Best,
>> >>>> Stefanos
>> >>>>
>> >>>> Στις Τρί, 10 Μαρ 2020 στις 1:04 π.μ., ο/η Hamilton Tobon
Mosquera via
>> >>>> llvm-dev
<llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>>
>> >>>> έγραψε:
>> >>>> Greetings,
>> >>>>
>> >>>>
>> >>>> I'm an LLVM newcomer interested in participating in the next GSoC
>> under
>> >>>> the project "Improve parallelism-aware analyses and
>> optimizations". I've
>> >>>> already read "Compiler Optimizations for OpenMP" and "Compiler
>> >>>> Optimizations for Parallel Programs" both by Doerfert and Finkel.
>> Also,
>> >>>> I've watched their two LLVM meeting developers conferences
>> "Representing
>> >>>> parallelism within LLVM" and "Optimizing indirections, using
>> >>>> abstractions without ...", so I have some background in the
problems
>> >>>> they are trying to tackle. Also, I'm close to solve one of
the bugs
>> >>>> posted in bugzilla, which has given me a broader idea of how the
>> OpenMP
>> >>>> pragmas are translated. The bus is this:
>> >>>>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&amp;data=01|01|htobonm%40eafit.edu.co|27650f5cf4e24c5920be08d7cf3c8209|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=UbADZ7YRMOi3%2FCVt%2FGvUMYCTw%2Bgmr8psk2QPVEOC8CE%3D&amp;reserved=0.
>> >>>>
>> >>>>
>> >>>> I tried contacting you Johannes but with your Argonne email,
I suppose
>> >>>> that's why you couldn't answer. Could you tell me more about the
>> >>>> project?, what variants do you have?, or what can I focus my
>> application
>> >>>> on?. Should I start with what you recommended to Emanuel?.
>> >>>>
>> >>>> Thank you.
>> >>>>
>> >>>> _______________________________________________
>> >>>> LLVM Developers mailing list
>> >>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >>>>
>>
llvm-dev Info Page
>> >>>>
>> >>>> La información contenida en este correo electrónico está dirigida
>> >>>> únicamente a su destinatario y puede contener información
>> >>>> confidencial, material privilegiado o información protegida por
>> >>>> derecho de autor. Está prohibida cualquier copia, utilización,
>> >>>> indebida retención, modificación, difusión, distribución o
>> >>>> reproducción total o parcial. Si usted recibe este mensaje
por error,
>> >>>> por favor contacte al remitente y elimínelo. La información aquí
>> >>>> contenida es responsabilidad exclusiva de su remitente por lo
tanto
>> >>>> la Universidad EAFIT no se hace responsable de lo que el mensaje
>> >>>> contenga. The information contained in this email is
addressed to its
>> >>>> recipient only and may contain confidential information,
privileged
>> >>>> material or information protected by copyright. Its
prohibited any
>> >>>> copy, use, improper retention, modification, dissemination,
>> >>>> distribution or total or partial reproduction. If you receive
this
>> >>>> message by error, please contact the sender and delete it. The
>> >>>> information contained herein is the sole responsibility of
the sender
>> >>>> therefore Universidad EAFIT is not responsible for what the
message
>> >>>> contains.
>> >>>> La información contenida en este correo electrónico está dirigida
>> >>>> únicamente a su destinatario y puede contener información
>> >>>> confidencial, material privilegiado o información protegida por
>> >>>> derecho de autor. Está prohibida cualquier copia, utilización,
>> >>>> indebida retención, modificación, difusión, distribución o
>> >>>> reproducción total o parcial. Si usted recibe este mensaje
por error,
>> >>>> por favor contacte al remitente y elimínelo. La información aquí
>> >>>> contenida es responsabilidad exclusiva de su remitente por lo
tanto
>> >>>> la Universidad EAFIT no se hace responsable de lo que el mensaje
>> >>>> contenga. The information contained in this email is
addressed to its
>> >>>> recipient only and may contain confidential information,
privileged
>> >>>> material or information protected by copyright. Its
prohibited any
>> >>>> copy, use, improper retention, modification, dissemination,
>> >>>> distribution or total or partial reproduction. If you receive
this
>> >>>> message by error, please contact the sender and delete it. The
>> >>>> information contained herein is the sole responsibility of
the sender
>> >>>> therefore Universidad EAFIT is not responsible for what the
message
>> >>>> contains.
>> >>>>
>> >>>>
>> >>>> _______________________________________________
>> >>>> LLVM Developers mailing list
>> >>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >>>>
>>
llvm-dev Info Page
>> >>>>
>> >>>>
>> >>>> La información contenida en este correo electrónico está dirigida
>> >>>> únicamente a su destinatario y puede contener información
>> >>>> confidencial, material privilegiado o información protegida por
>> >>>> derecho de autor. Está prohibida cualquier copia, utilización,
>> >>>> indebida retención, modificación, difusión, distribución o
>> >>>> reproducción total o parcial. Si usted recibe este mensaje
por error,
>> >>>> por favor contacte al remitente y elimínelo. La información aquí
>> >>>> contenida es responsabilidad exclusiva de su remitente por lo
tanto
>> >>>> la Universidad EAFIT no se hace responsable de lo que el mensaje
>> >>>> contenga. The information contained in this email is
addressed to its
>> >>>> recipient only and may contain confidential information,
privileged
>> >>>> material or information protected by copyright. Its
prohibited any
>> >>>> copy, use, improper retention, modification, dissemination,
>> >>>> distribution or total or partial reproduction. If you receive
this
>> >>>> message by error, please contact the sender and delete it. The
>> >>>> information contained herein is the sole responsibility of
the sender
>> >>>> therefore Universidad EAFIT is not responsible for what the
message
>> >>>> contains.
>> >>> _______________________________________________
>> >>> LLVM Developers mailing list
>> >>> llvm-dev@lists.llvm.org
>> >>>
>>
llvm-dev Info Page
>> >>>
>> >>> La información contenida en este correo electrónico está dirigida
>> >>> únicamente a su destinatario y puede contener información
>> >>> confidencial, material privilegiado o información protegida por
>> >>> derecho de autor. Está prohibida cualquier copia, utilización,
>> >>> indebida retención, modificación, difusión, distribución o
>> >>> reproducción total o parcial. Si usted recibe este mensaje por
error,
>> >>> por favor contacte al remitente y elimínelo. La información aquí
>> >>> contenida es responsabilidad exclusiva de su remitente por lo
tanto la
>> >>> Universidad EAFIT no se hace responsable de lo que el mensaje
>> >>> contenga. The information contained in this email is addressed
to its
>> >>> recipient only and may contain confidential information,
privileged
>> >>> material or information protected by copyright. Its prohibited any
>> >>> copy, use, improper retention, modification, dissemination,
>> >>> distribution or total or partial reproduction. If you receive this
>> >>> message by error, please contact the sender and delete it. The
>> >>> information contained herein is the sole responsibility of the
sender
>> >>> therefore Universidad EAFIT is not responsible for what the
message
>> >>> contains.
>> >> La información contenida en este correo electrónico está dirigida
>> únicamente a su destinatario y puede contener información confidencial,
>> material privilegiado o información protegida por derecho de autor. Está
>> prohibida cualquier copia, utilización, indebida retención,
>> modificación, difusión, distribución o reproducción total o parcial. Si
>> usted recibe este mensaje por error, por favor contacte al remitente y
>> elimínelo. La información aquí contenida es responsabilidad exclusiva de
>> su remitente por lo tanto la Universidad EAFIT no se hace responsable de
>> lo que el mensaje contenga. The information contained in this email is
>> addressed to its recipient only and may contain confidential
>> information, privileged material or information protected by copyright.
>> Its prohibited any copy, use, improper retention, modification,
>> dissemination, distribution or total or partial reproduction. If you
>> receive this message by error, please contact the sender and delete it.
>> The information contained herein is the sole responsibility of the
>> sender therefore Universidad EAFIT is not responsible for what the
>> message contains.
>> >
>>
>>
>> La información contenida en este correo electrónico está dirigida
únicamente a su destinatario y puede contener información confidencial,
material privilegiado o información protegida por derecho de autor. Está
prohibida cualquier copia, utilización, indebida retención,
modificación, difusión, distribución o reproducción total o parcial. Si
usted recibe este mensaje por error, por favor contacte al remitente y
elimínelo. La información aquí contenida es responsabilidad exclusiva de
su remitente por lo tanto la Universidad EAFIT no se hace responsable de
lo que el mensaje contenga. The information contained in this email is
addressed to its recipient only and may contain confidential
information, privileged material or information protected by copyright.
Its prohibited any copy, use, improper retention, modification,
dissemination, distribution or total or partial reproduction. If you
receive this message by error, please contact the sender and delete it.
The information contained herein is the sole responsibility of the
sender therefore Universidad EAFIT is not responsible for what the
message contains.

Hi Johannes,

Sorry about the failing tests. I completely forgot about running all of them, I only ran the tests for the openmp ipo transforms.

I found the GH issue in cache. I can post it here if you want.

I analyzed these runtime calls, __tgt_target_teams_[nowait]. What I don't understand about the problem we want to solve (splitting the memory transfer into issue and wait and legally moving code inside) is that the memory transfer is made inside those runtime calls. Therefore, the runtime calls to the memory transfer cannot be analyzed in the IR. Could you please explain a bit further if the memory transfer should split in the IR? or split __tgt_target_teams_[nowait] in such a way that the memory transfer is made in the first part, and the wait in the second part.

Thanks.

Hamilton.

Hi Johannes,

>
> Sorry about the failing tests. I completely forgot about running all
> of them, I only ran the tests for the openmp ipo transforms.

No worries, it happens.

> I found the GH issue in cache. I can post it here if you want.

Yes, please.

> I analyzed these runtime calls, __tgt_target_teams_[nowait]. What I
> don't understand about the problem we want to solve (splitting the
> memory transfer into issue and wait and legally moving code inside) is
> that the memory transfer is made inside those runtime calls.
> Therefore, the runtime calls to the memory transfer cannot be analyzed
> in the IR. Could you please explain a bit further if the memory
> transfer should split in the IR? or split __tgt_target_teams_[nowait]
> in such a way that the memory transfer is made in the first part, and
> the wait in the second part.

Exactly what you said in the last sentence. Split the runtime call into
two, issue and wait. Then we move issue "up" the execution path and wait
"down".

Cheers,
Johannes

> Thanks.
>
> Hamilton.
>
>>
>> > Hi Johannes,
>> >
>> > Yes, I'm highly interested, I already started. It would be nice
if you
>> > post the details again. I didn't take note of them, I was always
>> > looking at the issue in GH.
>>
>> I asked if we can make the readonly but unsure yet. I can write it again
>> but it'll take ma some time.
>>
>>
>> > The problem I'm facing right now is this:
>> > The first task is identifying OpenMP API calls referred to
device memory
>> > allocation. I first understood how the compiler driver
orchestrates the
>> > process of offloading OpenMP regions to devices, and now I'm
looking at
>> > the step of the pipeline that generates LLVM IR for both host and
>> > device. I wrote a simple C program that uses OpenMP to map a
local array
>> > to the device main memory. Then I'm looking at the device and
host LLVM
>> > IR representation of this program, but I can't find any runtime call
>> > referred to device memory allocation. In the host IR I see that the
>> > array is passed to the outlined function and stored in a
register inside
>> > the outlined function, but there's no memory allocation runtime
call.
>> > And in the device IR the array is passed to the main offloading
>> > function.
>> > Any suggestion or documentation on how that memory transfer is
>> > achieved?, and how to track its API calls. It might be that I'm not
>> > using the proper omp directive, so here is C code:
>>
>> Take a look at

, the mapping request,
>>
>> which may or may not cause data-transfers, are part of the call to
>> __tgt_target_teams[_nowait] and __tgt_target_data_update[_nowait]. Start
>> by reading the impl. of those two functions in the openmp runtime
>> (llvm-project/openmp/...).
>>
>> Let me know if this helps and once you have new questions.
>>
>> Cheers,
>> Johannes
>>
>> P.S. I merged your first patch. You didn't run all test before. I did
>> and noticed that we failed one because of type mismatches. Fixed them
>> before I merged your patch.
>>
>> >
>> > #include <omp.h>
>> > #include <stdio.h>
>> >
>> > #define size 1000000
>> > double a[size];
>> >
>> > int main() {
>> >
>> > #pragma omp target map(to: a[0:size])
>> > #pragma omp teams distribute
>> > for (int i = 0; i < size; ++i) {
>> > a[i] = a[i]*(i*i)/2;
>> > }
>> >
>> > return 0;
>> > }
>> >
>> >
>> > PD: The command to compile is this:
>> > clang -v -save-temps \
>> > -I/path/to/llvm/release/9.x/include -fopenmp \
>> > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
>> > -march=sm_61 memory_transfer.c -o memory_transfer
>> >
>> > Thanks.
>> > Hamilton.
>> >
>> >> > Ahhh I understand. So I believe the project you posted is still
>> >> opened. Do you any advice on the question I posted before?.
>> >>
>> >> Yes. Still open. You should use this as the main part of the GSoC
>> >> proposal (if you are interested). Can you repeat the question?
>> >>
>> >> >
>> >> > Thank you.
>> >> >
>> >> >> I did not. The issues sections on GH was removed.
>> >> >>
>> >> >> ________________________________________
>> >> >> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >> >> Sent: Sunday, March 22, 2020 12:21
>> >> >> To: llvm-dev@lists.llvm.org
>> >> >> Cc: Doerfert, Johannes
>> >> >> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >>
>> >> >> Hi there Johannes!,
>> >> >>
>> >> >> I was checking out the issue (about latency hiding when
>> transferring
>> >> >> memory from host to device) you posted about a week ago, but
>> it's not
>> >> >> there anymore. Did you remove it?.
>> >> >>
>> >> >> Thanks.
>> >> >>
>> >> >> Hamilton.
>> >> >>
>> >> >>> Hi there!,
>> >> >>>
>> >> >>> Right now I'm starting with the project proposed by
Johannes here:
>> >> >>>
>> >>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub.
>> >> >>>
>> >> >>>
>> >> >>> The first task is identifying OpenMP API calls referred
to device
>> >> memory
>> >> >>> allocation. I first understood how the compiler driver
>> orchestrates the
>> >> >>> process of offloading OpenMP regions to devices, and now I'm
>> looking at
>> >> >>> the step of the pipeline that generates LLVM IR for both
host and
>> >> >>> device. I wrote a simple C program that uses OpenMP to
map a local
>> >> array
>> >> >>> to the device main memory. Then I'm looking at the device and
>> host LLVM
>> >> >>> IR representation of this program, but I can't find any
>> runtime call
>> >> >>> referred to device memory allocation. In the host IR I see
>> that the
>> >> >>> array is passed to the outlined function and stored in a
register
>> >> inside
>> >> >>> the outlined function, but there's no memory allocation
>> runtime call.
>> >> >>> And in the device IR the array is passed to the main
offloading
>> >> >>> function. Any suggestion or documentation on how that memory
>> >> transfer is
>> >> >>> achieved?, and how to track its API calls. It might be that
>> I'm not
>> >> >>> using the proper omp directive, so here is C code:
>> >> >>>
>> >> >>>
>> >> >>> #include <omp.h>
>> >> >>> #include <stdio.h>
>> >> >>>
>> >> >>> #define size 1000000
>> >> >>> double a[size];
>> >> >>>
>> >> >>> int main() {
>> >> >>>
>> >> >>> #pragma omp target map(to: a[0:size])
>> >> >>> #pragma omp teams distribute
>> >> >>> for (int i = 0; i < size; ++i) {
>> >> >>> a[i] = a[i]*(i*i)/2;
>> >> >>> }
>> >> >>> return 0;
>> >> >>>
>> >> >>> PD: The command to compile is this: clang -v -save-temps
>> >> >>> -I/path/to/llvm/release/9.x/include -fopenmp
>> >> >>> -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target
-march=sm_61
>> >> >>> memory_transfer.c -o memory_transfer
>> >> >>>
>> >> >>> Thank you all.
>> >> >>>
>> >> >>> Hamilton.
>> >> >>>
>> >> >>>
>> >> >>>> Hi Hamilton,
>> >> >>>>
>> >> >>>> I personally don't believe the effort to make the IR
>> >> >>>> "parallelism-aware" is worth it right now.
>> >> >>>> What I propose is something else, see for example
>> >> >>>>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fyoutu.be%2FzfiHaPaoQPc&amp;data=01|01|htobonm%40eafit.edu.co|1c12e490f3804f6a772d08d7cf570a3a|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=PS4a99t8C7cFxrhlDsCma2nzGHGPdPne8hemx2lr%2Fms%3D&amp;reserved=0
>> >> >>>> and the OpenMPOpt pass.
>> >> >>>> There are multiple bigger tasks towards better
offloading, one of
>> >> >>>> which is described here
>> >> >>>>
>> >>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub
>> >> >>>> It might be interesting to add the transfer functions used
>> for memory
>> >> >>>> transfer to the OMPKinds.def file so they are known in the
>> >> >>>> OpenMPOpt.cpp pass.
>> >> >>>> Another small task would be to eliminate trivially redundant
>> ones,
>> >> >>>> that is if you have copy X from A to B directly followed by
>> copy X
>> >> >>>> from B to A, remove both.
>> >> >>>>
>> >> >>>> Does this make sense?
>> >> >>>>
>> >> >>>> Cheers,
>> >> >>>> Johannes
>> >> >>>>
>> >> >>>>
>> >> >>>>
>> >> >>>>
>> >> >>>>
>> >> >>>> ---------------------------------------
>> >> >>>> Johannes Doerfert
>> >> >>>> Researcher
>> >> >>>>
>> >> >>>> Argonne National Laboratory
>> >> >>>> Lemont, IL 60439, USA
>> >> >>>>
>> >> >>>> jdoerfert@anl.gov
>> >> >>>>
>> >> >>>> ________________________________________
>> >> >>>> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >> >>>> Sent: Wednesday, March 18, 2020 19:14
>> >> >>>> To: llvm-dev@lists.llvm.org
>> >> >>>> Cc: Doerfert, Johannes
>> >> >>>> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >>>>
>> >> >>>> Hi there!,
>> >> >>>>
>> >> >>>> I think I'm done with this patch: D76058. If that is
everything
>> >> >>>> concerning that, could you please assign me another task.
>> >> >>>>
>> >> >>>> I've been reading about Concurrent Control Flow Graphs,
which
>> would
>> >> >>>> be awesome for LLVM because it would be kind of a starting
>> point for
>> >> >>>> optimizations for parallel programs in general. The thing is
>> that it
>> >> >>>> seems more complicated than it is presented in the papers, I
>> think
>> >> >>>> that's why it's not implemented in LLVM yet. So, I think
I'll try
>> >> >>>> searching for something else. Nevertheless, don't you
think that
>> >> >>>> implementing this is worth the effort?. The papers I'm
>> reading are
>> >> >>>> from 1990s, I don't understand why they aren't widely
>> implemented in
>> >> >>>> big compilers xd.
>> >> >>>>
>> >> >>>> I saw this commit from you:
>> >> >>>>
>> >>
>>
⚙ D29250 [PIR] Add fork, join and halt instructions.
>> >> >>>> Is there something I can help with concerning this?.
>> >> >>>>
>> >> >>>> I would like to do contribute to the OpenMP target
offloading. Is
>> >> >>>> there anything I can help with?.
>> >> >>>>
>> >> >>>>
>> >> >>>> Hi there!,
>> >> >>>>
>> >> >>>> As part of my application process to the next GSoC I'm
>> working on the
>> >> >>>> TODO in OpenMPOpt.cpp line 437:
>> >> >>>>
>> >> >>>> // TODO: We should validate the declaration agains the
>> types we
>> >> >>>> expect.
>> >> >>>>
>> >> >>>> Link:
>> >> >>>>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fblob%2Fmaster%2Fllvm%2Flib%2FTransforms%2FIPO%2FOpenMPOpt.cpp%23L437&amp;data=01|01|htobonm%40eafit.edu.co|1c12e490f3804f6a772d08d7cf570a3a|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%3D&amp;reserved=0
>> >> >>>>
>> >> >>>> I have a question. When there is a mismatch in the types
>> (return type
>> >> >>>> or argument types) between the function declaration
found and the
>> >> >>>> runtime library function, what should I do?. Show an
error or a
>> >> >>>> warning?, continuing with the pass or exiting the program?.
>> >> >>>>
>> >> >>>> Thanks.
>> >> >>>>
>> >> >>>>
>> >> >>>> Hi Hamilton,
>> >> >>>>
>> >> >>>> Thanks for your interest in LLVM!
>> >> >>>>
>> >> >>>> IMHO you have a good start. I'll try to help by CC'ing
>> Johannes. Note
>> >> >>>> that it's difficult for LLVM developers
>> >> >>>> (including GSoC mentors and including Johannes) to keep up
>> with every
>> >> >>>> discussion on llvm-dev.
>> >> >>>> So, you can CC them to make their life easier. :slight_smile:
>> >> >>>>
>> >> >>>> Johannes, I see a lot of people interested in this project.
>> Maybe a
>> >> >>>> list of starting points, patches etc.
>> >> >>>> could help.
>> >> >>>>
>> >> >>>> Best,
>> >> >>>> Stefanos
>> >> >>>>
>> >> >>>> Στις Τρί, 10 Μαρ 2020 στις 1:04 π.μ., ο/η Hamilton Tobon
>> Mosquera via
>> >> >>>> llvm-dev
>> <llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>>
>> >> >>>> έγραψε:
>> >> >>>> Greetings,
>> >> >>>>
>> >> >>>>
>> >> >>>> I'm an LLVM newcomer interested in participating in the
next GSoC
>> >> under
>> >> >>>> the project "Improve parallelism-aware analyses and
>> >> optimizations". I've
>> >> >>>> already read "Compiler Optimizations for OpenMP" and
"Compiler
>> >> >>>> Optimizations for Parallel Programs" both by Doerfert
and Finkel.
>> >> Also,
>> >> >>>> I've watched their two LLVM meeting developers conferences
>> >> "Representing
>> >> >>>> parallelism within LLVM" and "Optimizing indirections, using
>> >> >>>> abstractions without ...", so I have some background in the
>> problems
>> >> >>>> they are trying to tackle. Also, I'm close to solve one of
>> the bugs
>> >> >>>> posted in bugzilla, which has given me a broader idea of
how the
>> >> OpenMP
>> >> >>>> pragmas are translated. The bus is this:
>> >> >>>>
>> >>
>>
40253 – combined target teams implements shared clause as firstprivate.
>> >> >>>>
>> >> >>>>
>> >> >>>> I tried contacting you Johannes but with your Argonne email,
>> I suppose
>> >> >>>> that's why you couldn't answer. Could you tell me more
about the
>> >> >>>> project?, what variants do you have?, or what can I focus my
>> >> application
>> >> >>>> on?. Should I start with what you recommended to Emanuel?.
>> >> >>>>
>> >> >>>> Thank you.
>> >> >>>>
>> >> >>>> _______________________________________________
>> >> >>>> LLVM Developers mailing list
>> >> >>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >> >>>>
>> >>
>>
llvm-dev Info Page
>> >> >>>>
>> >> >>>> La información contenida en este correo electrónico está
dirigida
>> >> >>>> únicamente a su destinatario y puede contener información
>> >> >>>> confidencial, material privilegiado o información
protegida por
>> >> >>>> derecho de autor. Está prohibida cualquier copia,
utilización,
>> >> >>>> indebida retención, modificación, difusión, distribución o
>> >> >>>> reproducción total o parcial. Si usted recibe este mensaje
>> por error,
>> >> >>>> por favor contacte al remitente y elimínelo. La
información aquí
>> >> >>>> contenida es responsabilidad exclusiva de su remitente
por lo
>> tanto
>> >> >>>> la Universidad EAFIT no se hace responsable de lo que el
mensaje
>> >> >>>> contenga. The information contained in this email is
>> addressed to its
>> >> >>>> recipient only and may contain confidential information,
>> privileged
>> >> >>>> material or information protected by copyright. Its
>> prohibited any
>> >> >>>> copy, use, improper retention, modification, dissemination,
>> >> >>>> distribution or total or partial reproduction. If you
receive
>> this
>> >> >>>> message by error, please contact the sender and delete
it. The
>> >> >>>> information contained herein is the sole responsibility of
>> the sender
>> >> >>>> therefore Universidad EAFIT is not responsible for what the
>> message
>> >> >>>> contains.
>> >> >>>> La información contenida en este correo electrónico está
dirigida
>> >> >>>> únicamente a su destinatario y puede contener información
>> >> >>>> confidencial, material privilegiado o información
protegida por
>> >> >>>> derecho de autor. Está prohibida cualquier copia,
utilización,
>> >> >>>> indebida retención, modificación, difusión, distribución o
>> >> >>>> reproducción total o parcial. Si usted recibe este mensaje
>> por error,
>> >> >>>> por favor contacte al remitente y elimínelo. La
información aquí
>> >> >>>> contenida es responsabilidad exclusiva de su remitente
por lo
>> tanto
>> >> >>>> la Universidad EAFIT no se hace responsable de lo que el
mensaje
>> >> >>>> contenga. The information contained in this email is
>> addressed to its
>> >> >>>> recipient only and may contain confidential information,
>> privileged
>> >> >>>> material or information protected by copyright. Its
>> prohibited any
>> >> >>>> copy, use, improper retention, modification, dissemination,
>> >> >>>> distribution or total or partial reproduction. If you
receive
>> this
>> >> >>>> message by error, please contact the sender and delete
it. The
>> >> >>>> information contained herein is the sole responsibility of
>> the sender
>> >> >>>> therefore Universidad EAFIT is not responsible for what the
>> message
>> >> >>>> contains.
>> >> >>>>
>> >> >>>>
>> >> >>>> _______________________________________________
>> >> >>>> LLVM Developers mailing list
>> >> >>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >> >>>>
>> >>
>>
llvm-dev Info Page
>> >> >>>>
>> >> >>>>
>> >> >>>> La información contenida en este correo electrónico está
dirigida
>> >> >>>> únicamente a su destinatario y puede contener información
>> >> >>>> confidencial, material privilegiado o información
protegida por
>> >> >>>> derecho de autor. Está prohibida cualquier copia,
utilización,
>> >> >>>> indebida retención, modificación, difusión, distribución o
>> >> >>>> reproducción total o parcial. Si usted recibe este mensaje
>> por error,
>> >> >>>> por favor contacte al remitente y elimínelo. La
información aquí
>> >> >>>> contenida es responsabilidad exclusiva de su remitente
por lo
>> tanto
>> >> >>>> la Universidad EAFIT no se hace responsable de lo que el
mensaje
>> >> >>>> contenga. The information contained in this email is
>> addressed to its
>> >> >>>> recipient only and may contain confidential information,
>> privileged
>> >> >>>> material or information protected by copyright. Its
>> prohibited any
>> >> >>>> copy, use, improper retention, modification, dissemination,
>> >> >>>> distribution or total or partial reproduction. If you
receive
>> this
>> >> >>>> message by error, please contact the sender and delete
it. The
>> >> >>>> information contained herein is the sole responsibility of
>> the sender
>> >> >>>> therefore Universidad EAFIT is not responsible for what the
>> message
>> >> >>>> contains.
>> >> >>> _______________________________________________
>> >> >>> LLVM Developers mailing list
>> >> >>> llvm-dev@lists.llvm.org
>> >> >>>
>> >>
>>
llvm-dev Info Page
>> >> >>>
>> >> >>> La información contenida en este correo electrónico está
dirigida
>> >> >>> únicamente a su destinatario y puede contener información
>> >> >>> confidencial, material privilegiado o información
protegida por
>> >> >>> derecho de autor. Está prohibida cualquier copia,
utilización,
>> >> >>> indebida retención, modificación, difusión, distribución o
>> >> >>> reproducción total o parcial. Si usted recibe este
mensaje por
>> error,
>> >> >>> por favor contacte al remitente y elimínelo. La
información aquí
>> >> >>> contenida es responsabilidad exclusiva de su remitente por lo
>> tanto la
>> >> >>> Universidad EAFIT no se hace responsable de lo que el mensaje
>> >> >>> contenga. The information contained in this email is
addressed
>> to its
>> >> >>> recipient only and may contain confidential information,
>> privileged
>> >> >>> material or information protected by copyright. Its
prohibited any
>> >> >>> copy, use, improper retention, modification, dissemination,
>> >> >>> distribution or total or partial reproduction. If you
receive this
>> >> >>> message by error, please contact the sender and delete
it. The
>> >> >>> information contained herein is the sole responsibility
of the
>> sender
>> >> >>> therefore Universidad EAFIT is not responsible for what the
>> message
>> >> >>> contains.
>> >> >> La información contenida en este correo electrónico está
dirigida
>> >> únicamente a su destinatario y puede contener información
confidencial,
>> >> material privilegiado o información protegida por derecho de
autor. Está
>> >> prohibida cualquier copia, utilización, indebida retención,
>> >> modificación, difusión, distribución o reproducción total o
parcial. Si
>> >> usted recibe este mensaje por error, por favor contacte al
remitente y
>> >> elimínelo. La información aquí contenida es responsabilidad
exclusiva de
>> >> su remitente por lo tanto la Universidad EAFIT no se hace
responsable de
>> >> lo que el mensaje contenga. The information contained in this
email is
>> >> addressed to its recipient only and may contain confidential
>> >> information, privileged material or information protected by
copyright.
>> >> Its prohibited any copy, use, improper retention, modification,
>> >> dissemination, distribution or total or partial reproduction.
If you
>> >> receive this message by error, please contact the sender and
delete it.
>> >> The information contained herein is the sole responsibility of the
>> >> sender therefore Universidad EAFIT is not responsible for what the
>> >> message contains.
>> >> >
>> >>
>> >>
>> >> La información contenida en este correo electrónico está dirigida
>> únicamente a su destinatario y puede contener información confidencial,
>> material privilegiado o información protegida por derecho de autor. Está
>> prohibida cualquier copia, utilización, indebida retención,
>> modificación, difusión, distribución o reproducción total o parcial. Si
>> usted recibe este mensaje por error, por favor contacte al remitente y
>> elimínelo. La información aquí contenida es responsabilidad exclusiva de
>> su remitente por lo tanto la Universidad EAFIT no se hace responsable de
>> lo que el mensaje contenga. The information contained in this email is
>> addressed to its recipient only and may contain confidential
>> information, privileged material or information protected by copyright.
>> Its prohibited any copy, use, improper retention, modification,
>> dissemination, distribution or total or partial reproduction. If you
>> receive this message by error, please contact the sender and delete it.
>> The information contained herein is the sole responsibility of the
>> sender therefore Universidad EAFIT is not responsible for what the
>> message contains.
>> >
>>
>>
>> La información contenida en este correo electrónico está dirigida
únicamente a su destinatario y puede contener información confidencial,
material privilegiado o información protegida por derecho de autor. Está
prohibida cualquier copia, utilización, indebida retención,
modificación, difusión, distribución o reproducción total o parcial. Si
usted recibe este mensaje por error, por favor contacte al remitente y
elimínelo. La información aquí contenida es responsabilidad exclusiva de
su remitente por lo tanto la Universidad EAFIT no se hace responsable de
lo que el mensaje contenga. The information contained in this email is
addressed to its recipient only and may contain confidential
information, privileged material or information protected by copyright.
Its prohibited any copy, use, improper retention, modification,
dissemination, distribution or total or partial reproduction. If you
receive this message by error, please contact the sender and delete it.
The information contained herein is the sole responsibility of the
sender therefore Universidad EAFIT is not responsible for what the
message contains.

Hi Johannes,

I've been thinking on how to implement a solution for the problem and I have some questions:

1. You mention in the problem description that the current asynchronous functions can be used, and probably modified. The problem is that those runtime functions with _nowait at the end are not asynchronous. Looking at their implementations, they just wait for other tasks is there are pending tasks, and then call its synchronous versions. So, is it a good idea to wrap them in functions that call them as asynchronous (using std::async), then return an object to wait on?. Then, waiting on the object returned.

2. I'm thinking on replacing the data transfer inside __tgt_target_teams with __tgt_target_data_begin. The problem with this is that __tgt_target_data_begin does not initialize [first]private variables, and doesn't check for lambda mappings. Therefore, my idea to go around this problem is to create another wrapper function F, that calls __tgt_target_data_begin, and executes the necessary code for lambda mappings check and [first]private variables, and finally, calling F as the asynchronous function. What is good about this, is that the code concerning data transfer inside __tgt_target_teams can be replaced with F.

3. Also, I was wondering if I must to be aware of (split them into issue and wait) all the other constructs that imply data transfers (update, data enter/exit), or just focus on __tgt_target[_nowait] and __tgt_target_teams[_nowait].

4. This might be a dumb question, but just to make sure. Do I have to move the wait call into the offloaded region just before the data is used (because I think this would imply analyzing the offloaded bitcode, which is in another module)?, or move it just before the kernel is executed?.

5. Do you have a template for the GSoC proposal. That is, is it plain text or pdf or ...?. Also, do you have sort of a structure that we have to follow?, for example, Abstract -> Benefits -> Deliverables ..., or do we have freedom on that?. Your post in hte GSoC page only mentions that we must include our background, but nothing about the format of the proposal.

Thank you.

Hamilton.

On``` 25/03/20 11:41 a. m., Doerfert, Johannes wrote:

Hi Johannes,

>
> I've been thinking on how to implement a solution for the problem and
> I have some questions:
>
> 1. You mention in the problem description that the current
> asynchronous functions can be used, and probably modified. The problem
> is that those runtime functions with _nowait at the end are not
> asynchronous. Looking at their implementations, they just wait for
> other tasks is there are pending tasks, and then call its synchronous
> versions. So, is it a good idea to wrap them in functions that call
> them as asynchronous (using std::async), then return an object to wait
> on?. Then, waiting on the object returned.

Take a look at D77005. Shilei introduced some real askync versions for a
different reason already. You can build on top of that patch and extend
it if needed further.

> 2. I'm thinking on replacing the data transfer inside
> __tgt_target_teams with __tgt_target_data_begin. The problem with this
> is that __tgt_target_data_begin does not initialize [first]private
> variables, and doesn't check for lambda mappings. Therefore, my idea
> to go around this problem is to create another wrapper function F,
> that calls __tgt_target_data_begin, and executes the necessary code
> for lambda mappings check and [first]private variables, and finally,
> calling F as the asynchronous function. What is good about this, is
> that the code concerning data transfer inside __tgt_target_teams can
> be replaced with F.

Interesting. I'll need to see small examples. For now, focus on the
simple case, we'll get to the rest later. However, we should really
start by creating test cases that show the initial code and FIXMEs that
describe how we want it to look like.

> 3. Also, I was wondering if I must to be aware of (split them into
> issue and wait) all the other constructs that imply data transfers
> (update, data enter/exit), or just focus on __tgt_target[_nowait] and
> __tgt_target_teams[_nowait].

We want to eventually split all transfers into two but let's start with
a single one.

> 4. This might be a dumb question, but just to make sure. Do I have to
> move the wait call into the offloaded region just before the data is
> used (because I think this would imply analyzing the offloaded
> bitcode, which is in another module)?, or move it just before the
> kernel is executed?.

We split, then we move them apart. The cannot be moved over instructions
that might change behavior if we move them, e.g. an aliasing memory
operation or a kernel call that uses the data.

> 5. Do you have a template for the GSoC proposal. That is, is it plain
> text or pdf or ...?. Also, do you have sort of a structure that we
> have to follow?, for example, Abstract -> Benefits -> Deliverables
> ..., or do we have freedom on that?. Your post in hte GSoC page only
> mentions that we must include our background, but nothing about the
> format of the proposal.

I do not have one. I thought they might or they have an FAQ that
describes what to do. I know they have some required things, so check
their webpage. I think you have to use google docs, I might be wrong.
Your categories look good, make the deliverables not too specific but
describe more ideas to work on. We'll make tasks as we go as you are not
the only one working on this general area.

Feel free to send me a draft of the proposal for feedback.

Cheers,
Johannes

> Thank you.
>
> Hamilton.
>
>
> On``` 25/03/20 11:41 a. m., Doerfert, Johannes wrote:
>> > Hi Johannes,
>> >
>> > Sorry about the failing tests. I completely forgot about running all
>> > of them, I only ran the tests for the openmp ipo transforms.
>>
>> No worries, it happens.
>>
>>
>> > I found the GH issue in cache. I can post it here if you want.
>>
>> Yes, please.
>>
>>
>> > I analyzed these runtime calls, __tgt_target_teams_[nowait]. What I
>> > don't understand about the problem we want to solve (splitting the
>> > memory transfer into issue and wait and legally moving code
inside) is
>> > that the memory transfer is made inside those runtime calls.
>> > Therefore, the runtime calls to the memory transfer cannot be
analyzed
>> > in the IR. Could you please explain a bit further if the memory
>> > transfer should split in the IR? or split
__tgt_target_teams_[nowait]
>> > in such a way that the memory transfer is made in the first
part, and
>> > the wait in the second part.
>>
>> Exactly what you said in the last sentence. Split the runtime call into
>> two, issue and wait. Then we move issue "up" the execution path and wait
>> "down".
>>
>> Cheers,
>> Johannes
>>
>>
>> > Thanks.
>> >
>> > Hamilton.
>> >
>> >>
>> >> > Hi Johannes,
>> >> >
>> >> > Yes, I'm highly interested, I already started. It would be nice
>> if you
>> >> > post the details again. I didn't take note of them, I was
always
>> >> > looking at the issue in GH.
>> >>
>> >> I asked if we can make the readonly but unsure yet. I can write
it again
>> >> but it'll take ma some time.
>> >>
>> >>
>> >> > The problem I'm facing right now is this:
>> >> > The first task is identifying OpenMP API calls referred to
>> device memory
>> >> > allocation. I first understood how the compiler driver
>> orchestrates the
>> >> > process of offloading OpenMP regions to devices, and now I'm
>> looking at
>> >> > the step of the pipeline that generates LLVM IR for both
host and
>> >> > device. I wrote a simple C program that uses OpenMP to map a
>> local array
>> >> > to the device main memory. Then I'm looking at the device and
>> host LLVM
>> >> > IR representation of this program, but I can't find any
runtime call
>> >> > referred to device memory allocation. In the host IR I see
that the
>> >> > array is passed to the outlined function and stored in a
>> register inside
>> >> > the outlined function, but there's no memory allocation runtime
>> call.
>> >> > And in the device IR the array is passed to the main offloading
>> >> > function.
>> >> > Any suggestion or documentation on how that memory transfer is
>> >> > achieved?, and how to track its API calls. It might be that
I'm not
>> >> > using the proper omp directive, so here is C code:
>> >>
>> >> Take a look at
>>

>> , the mapping request,
>> >>
>> >> which may or may not cause data-transfers, are part of the call to
>> >> __tgt_target_teams[_nowait] and
__tgt_target_data_update[_nowait]. Start
>> >> by reading the impl. of those two functions in the openmp runtime
>> >> (llvm-project/openmp/...).
>> >>
>> >> Let me know if this helps and once you have new questions.
>> >>
>> >> Cheers,
>> >> Johannes
>> >>
>> >> P.S. I merged your first patch. You didn't run all test before.
I did
>> >> and noticed that we failed one because of type mismatches.
Fixed them
>> >> before I merged your patch.
>> >>
>> >> >
>> >> > #include <omp.h>
>> >> > #include <stdio.h>
>> >> >
>> >> > #define size 1000000
>> >> > double a[size];
>> >> >
>> >> > int main() {
>> >> >
>> >> > #pragma omp target map(to: a[0:size])
>> >> > #pragma omp teams distribute
>> >> > for (int i = 0; i < size; ++i) {
>> >> > a[i] = a[i]*(i*i)/2;
>> >> > }
>> >> >
>> >> > return 0;
>> >> > }
>> >> >
>> >> >
>> >> > PD: The command to compile is this:
>> >> > clang -v -save-temps \
>> >> > -I/path/to/llvm/release/9.x/include -fopenmp \
>> >> > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
>> >> > -march=sm_61 memory_transfer.c -o memory_transfer
>> >> >
>> >> > Thanks.
>> >> > Hamilton.
>> >> >
>> >> >> > Ahhh I understand. So I believe the project you posted
is still
>> >> >> opened. Do you any advice on the question I posted before?.
>> >> >>
>> >> >> Yes. Still open. You should use this as the main part of
the GSoC
>> >> >> proposal (if you are interested). Can you repeat the question?
>> >> >>
>> >> >> >
>> >> >> > Thank you.
>> >> >> >
>> >> >> >> I did not. The issues sections on GH was removed.
>> >> >> >>
>> >> >> >> ________________________________________
>> >> >> >> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >> >> >> Sent: Sunday, March 22, 2020 12:21
>> >> >> >> To: llvm-dev@lists.llvm.org
>> >> >> >> Cc: Doerfert, Johannes
>> >> >> >> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >> >>
>> >> >> >> Hi there Johannes!,
>> >> >> >>
>> >> >> >> I was checking out the issue (about latency hiding when
>> >> transferring
>> >> >> >> memory from host to device) you posted about a week
ago, but
>> >> it's not
>> >> >> >> there anymore. Did you remove it?.
>> >> >> >>
>> >> >> >> Thanks.
>> >> >> >>
>> >> >> >> Hamilton.
>> >> >> >>
>> >> >> >>> Hi there!,
>> >> >> >>>
>> >> >> >>> Right now I'm starting with the project proposed by
>> Johannes here:
>> >> >> >>>
>> >> >>
>> >>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub.
>> >> >> >>>
>> >> >> >>>
>> >> >> >>> The first task is identifying OpenMP API calls referred
>> to device
>> >> >> memory
>> >> >> >>> allocation. I first understood how the compiler driver
>> >> orchestrates the
>> >> >> >>> process of offloading OpenMP regions to devices, and
now I'm
>> >> looking at
>> >> >> >>> the step of the pipeline that generates LLVM IR for both
>> host and
>> >> >> >>> device. I wrote a simple C program that uses OpenMP to
>> map a local
>> >> >> array
>> >> >> >>> to the device main memory. Then I'm looking at the
device and
>> >> host LLVM
>> >> >> >>> IR representation of this program, but I can't find any
>> >> runtime call
>> >> >> >>> referred to device memory allocation. In the host IR
I see
>> >> that the
>> >> >> >>> array is passed to the outlined function and stored in a
>> register
>> >> >> inside
>> >> >> >>> the outlined function, but there's no memory allocation
>> >> runtime call.
>> >> >> >>> And in the device IR the array is passed to the main
>> offloading
>> >> >> >>> function. Any suggestion or documentation on how
that memory
>> >> >> transfer is
>> >> >> >>> achieved?, and how to track its API calls. It might
be that
>> >> I'm not
>> >> >> >>> using the proper omp directive, so here is C code:
>> >> >> >>>
>> >> >> >>>
>> >> >> >>> #include <omp.h>
>> >> >> >>> #include <stdio.h>
>> >> >> >>>
>> >> >> >>> #define size 1000000
>> >> >> >>> double a[size];
>> >> >> >>>
>> >> >> >>> int main() {
>> >> >> >>>
>> >> >> >>> #pragma omp target map(to: a[0:size])
>> >> >> >>> #pragma omp teams distribute
>> >> >> >>> for (int i = 0; i < size; ++i) {
>> >> >> >>> a[i] = a[i]*(i*i)/2;
>> >> >> >>> }
>> >> >> >>> return 0;
>> >> >> >>>
>> >> >> >>> PD: The command to compile is this: clang -v -save-temps
>> >> >> >>> -I/path/to/llvm/release/9.x/include -fopenmp
>> >> >> >>> -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target
>> -march=sm_61
>> >> >> >>> memory_transfer.c -o memory_transfer
>> >> >> >>>
>> >> >> >>> Thank you all.
>> >> >> >>>
>> >> >> >>> Hamilton.
>> >> >> >>>
>> >> >> >>>
>> >> >> >>>> Hi Hamilton,
>> >> >> >>>>
>> >> >> >>>> I personally don't believe the effort to make the IR
>> >> >> >>>> "parallelism-aware" is worth it right now.
>> >> >> >>>> What I propose is something else, see for example
>> >> >> >>>>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fyoutu.be%2FzfiHaPaoQPc&amp;data=01|01|htobonm%40eafit.edu.co|d2a1829610714bf161b508d7d0d2f5ae|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=zQQajI2aySpFNO4EieIW8u5I2ei5i6myfdescuN0%2FtY%3D&amp;reserved=0
>> >> >> >>>> and the OpenMPOpt pass.
>> >> >> >>>> There are multiple bigger tasks towards better
>> offloading, one of
>> >> >> >>>> which is described here
>> >> >> >>>>
>> >> >>
>> >>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub
>> >> >> >>>> It might be interesting to add the transfer
functions used
>> >> for memory
>> >> >> >>>> transfer to the OMPKinds.def file so they are known
in the
>> >> >> >>>> OpenMPOpt.cpp pass.
>> >> >> >>>> Another small task would be to eliminate trivially
redundant
>> >> ones,
>> >> >> >>>> that is if you have copy X from A to B directly
followed by
>> >> copy X
>> >> >> >>>> from B to A, remove both.
>> >> >> >>>>
>> >> >> >>>> Does this make sense?
>> >> >> >>>>
>> >> >> >>>> Cheers,
>> >> >> >>>> Johannes
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> ---------------------------------------
>> >> >> >>>> Johannes Doerfert
>> >> >> >>>> Researcher
>> >> >> >>>>
>> >> >> >>>> Argonne National Laboratory
>> >> >> >>>> Lemont, IL 60439, USA
>> >> >> >>>>
>> >> >> >>>> jdoerfert@anl.gov
>> >> >> >>>>
>> >> >> >>>> ________________________________________
>> >> >> >>>> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >> >> >>>> Sent: Wednesday, March 18, 2020 19:14
>> >> >> >>>> To: llvm-dev@lists.llvm.org
>> >> >> >>>> Cc: Doerfert, Johannes
>> >> >> >>>> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >> >>>>
>> >> >> >>>> Hi there!,
>> >> >> >>>>
>> >> >> >>>> I think I'm done with this patch: D76058. If that is
>> everything
>> >> >> >>>> concerning that, could you please assign me another
task.
>> >> >> >>>>
>> >> >> >>>> I've been reading about Concurrent Control Flow Graphs,
>> which
>> >> would
>> >> >> >>>> be awesome for LLVM because it would be kind of a
starting
>> >> point for
>> >> >> >>>> optimizations for parallel programs in general. The
thing is
>> >> that it
>> >> >> >>>> seems more complicated than it is presented in the
papers, I
>> >> think
>> >> >> >>>> that's why it's not implemented in LLVM yet. So, I
think
>> I'll try
>> >> >> >>>> searching for something else. Nevertheless, don't you
>> think that
>> >> >> >>>> implementing this is worth the effort?. The papers I'm
>> >> reading are
>> >> >> >>>> from 1990s, I don't understand why they aren't widely
>> >> implemented in
>> >> >> >>>> big compilers xd.
>> >> >> >>>>
>> >> >> >>>> I saw this commit from you:
>> >> >> >>>>
>> >> >>
>> >>
>>
⚙ D29250 [PIR] Add fork, join and halt instructions.
>> >> >> >>>> Is there something I can help with concerning this?.
>> >> >> >>>>
>> >> >> >>>> I would like to do contribute to the OpenMP target
>> offloading. Is
>> >> >> >>>> there anything I can help with?.
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> Hi there!,
>> >> >> >>>>
>> >> >> >>>> As part of my application process to the next GSoC I'm
>> >> working on the
>> >> >> >>>> TODO in OpenMPOpt.cpp line 437:
>> >> >> >>>>
>> >> >> >>>> // TODO: We should validate the declaration
agains the
>> >> types we
>> >> >> >>>> expect.
>> >> >> >>>>
>> >> >> >>>> Link:
>> >> >> >>>>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fblob%2Fmaster%2Fllvm%2Flib%2FTransforms%2FIPO%2FOpenMPOpt.cpp%23L437&amp;data=01|01|htobonm%40eafit.edu.co|d2a1829610714bf161b508d7d0d2f5ae|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=47sDHCNwfoSt%2B7HDCHg5xe7yTK2IYYPOY6k2NxvGaj8%3D&amp;reserved=0
>> >> >> >>>>
>> >> >> >>>> I have a question. When there is a mismatch in the
types
>> >> (return type
>> >> >> >>>> or argument types) between the function declaration
>> found and the
>> >> >> >>>> runtime library function, what should I do?. Show an
>> error or a
>> >> >> >>>> warning?, continuing with the pass or exiting the
program?.
>> >> >> >>>>
>> >> >> >>>> Thanks.
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> Hi Hamilton,
>> >> >> >>>>
>> >> >> >>>> Thanks for your interest in LLVM!
>> >> >> >>>>
>> >> >> >>>> IMHO you have a good start. I'll try to help by CC'ing
>> >> Johannes. Note
>> >> >> >>>> that it's difficult for LLVM developers
>> >> >> >>>> (including GSoC mentors and including Johannes) to
keep up
>> >> with every
>> >> >> >>>> discussion on llvm-dev.
>> >> >> >>>> So, you can CC them to make their life easier. :slight_smile:
>> >> >> >>>>
>> >> >> >>>> Johannes, I see a lot of people interested in this
project.
>> >> Maybe a
>> >> >> >>>> list of starting points, patches etc.
>> >> >> >>>> could help.
>> >> >> >>>>
>> >> >> >>>> Best,
>> >> >> >>>> Stefanos
>> >> >> >>>>
>> >> >> >>>> Στις Τρί, 10 Μαρ 2020 στις 1:04 π.μ., ο/η Hamilton
Tobon
>> >> Mosquera via
>> >> >> >>>> llvm-dev
>> >> <llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>>
>> >> >> >>>> έγραψε:
>> >> >> >>>> Greetings,
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> I'm an LLVM newcomer interested in participating in the
>> next GSoC
>> >> >> under
>> >> >> >>>> the project "Improve parallelism-aware analyses and
>> >> >> optimizations". I've
>> >> >> >>>> already read "Compiler Optimizations for OpenMP" and
>> "Compiler
>> >> >> >>>> Optimizations for Parallel Programs" both by Doerfert
>> and Finkel.
>> >> >> Also,
>> >> >> >>>> I've watched their two LLVM meeting developers
conferences
>> >> >> "Representing
>> >> >> >>>> parallelism within LLVM" and "Optimizing
indirections, using
>> >> >> >>>> abstractions without ...", so I have some
background in the
>> >> problems
>> >> >> >>>> they are trying to tackle. Also, I'm close to solve
one of
>> >> the bugs
>> >> >> >>>> posted in bugzilla, which has given me a broader
idea of
>> how the
>> >> >> OpenMP
>> >> >> >>>> pragmas are translated. The bus is this:
>> >> >> >>>>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&amp;data=01|01|htobonm%40eafit.edu.co|d2a1829610714bf161b508d7d0d2f5ae|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=r2Gq2MnAyC8jZzMYPHkEKRnL7p2teev2K3t9%2FdtE9es%3D&amp;reserved=0.
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> I tried contacting you Johannes but with your
Argonne email,
>> >> I suppose
>> >> >> >>>> that's why you couldn't answer. Could you tell me more
>> about the
>> >> >> >>>> project?, what variants do you have?, or what can I
focus my
>> >> >> application
>> >> >> >>>> on?. Should I start with what you recommended to
Emanuel?.
>> >> >> >>>>
>> >> >> >>>> Thank you.
>> >> >> >>>>
>> >> >> >>>> _______________________________________________
>> >> >> >>>> LLVM Developers mailing list
>> >> >> >>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >> >> >>>>
>> >> >>
>> >>
>>
llvm-dev Info Page
>> >> >> >>>>
>> >> >> >>>> La información contenida en este correo electrónico
está
>> dirigida
>> >> >> >>>> únicamente a su destinatario y puede contener
información
>> >> >> >>>> confidencial, material privilegiado o información
>> protegida por
>> >> >> >>>> derecho de autor. Está prohibida cualquier copia,
>> utilización,
>> >> >> >>>> indebida retención, modificación, difusión,
distribución o
>> >> >> >>>> reproducción total o parcial. Si usted recibe este
mensaje
>> >> por error,
>> >> >> >>>> por favor contacte al remitente y elimínelo. La
>> información aquí
>> >> >> >>>> contenida es responsabilidad exclusiva de su remitente
>> por lo
>> >> tanto
>> >> >> >>>> la Universidad EAFIT no se hace responsable de lo
que el
>> mensaje
>> >> >> >>>> contenga. The information contained in this email is
>> >> addressed to its
>> >> >> >>>> recipient only and may contain confidential
information,
>> >> privileged
>> >> >> >>>> material or information protected by copyright. Its
>> >> prohibited any
>> >> >> >>>> copy, use, improper retention, modification,
dissemination,
>> >> >> >>>> distribution or total or partial reproduction. If you
>> receive
>> >> this
>> >>�� >> >>>> message by error, please contact the sender and delete
>> it. The
>> >> >> >>>> information contained herein is the sole
responsibility of
>> >> the sender
>> >> >> >>>> therefore Universidad EAFIT is not responsible for
what the
>> >> message
>> >> >> >>>> contains.
>> >> >> >>>> La información contenida en este correo electrónico
está
>> dirigida
>> >> >> >>>> únicamente a su destinatario y puede contener
información
>> >> >> >>>> confidencial, material privilegiado o información
>> protegida por
>> >> >> >>>> derecho de autor. Está prohibida cualquier copia,
>> utilización,
>> >> >> >>>> indebida retención, modificación, difusión,
distribución o
>> >> >> >>>> reproducción total o parcial. Si usted recibe este
mensaje
>> >> por error,
>> >> >> >>>> por favor contacte al remitente y elimínelo. La
>> información aquí
>> >> >> >>>> contenida es responsabilidad exclusiva de su remitente
>> por lo
>> >> tanto
>> >> >> >>>> la Universidad EAFIT no se hace responsable de lo
que el
>> mensaje
>> >> >> >>>> contenga. The information contained in this email is
>> >> addressed to its
>> >> >> >>>> recipient only and may contain confidential
information,
>> >> privileged
>> >> >> >>>> material or information protected by copyright. Its
>> >> prohibited any
>> >> >> >>>> copy, use, improper retention, modification,
dissemination,
>> >> >> >>>> distribution or total or partial reproduction. If you
>> receive
>> >> this
>> >> >> >>>> message by error, please contact the sender and delete
>> it. The
>> >> >> >>>> information contained herein is the sole
responsibility of
>> >> the sender
>> >> >> >>>> therefore Universidad EAFIT is not responsible for
what the
>> >> message
>> >> >> >>>> contains.
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> _______________________________________________
>> >> >> >>>> LLVM Developers mailing list
>> >> >> >>>> llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >> >> >>>>
>> >> >>
>> >>
>>
llvm-dev Info Page
>> >> >> >>>>
>> >> >> >>>>
>> >> >> >>>> La información contenida en este correo electrónico
está
>> dirigida
>> >> >> >>>> únicamente a su destinatario y puede contener
información
>> >> >> >>>> confidencial, material privilegiado o información
>> protegida por
>> >> >> >>>> derecho de autor. Está prohibida cualquier copia,
>> utilización,
>> >> >> >>>> indebida retención, modificación, difusión,
distribución o
>> >> >> >>>> reproducción total o parcial. Si usted recibe este
mensaje
>> >> por error,
>> >> >> >>>> por favor contacte al remitente y elimínelo. La
>> información aquí
>> >> >> >>>> contenida es responsabilidad exclusiva de su remitente
>> por lo
>> >> tanto
>> >> >> >>>> la Universidad EAFIT no se hace responsable de lo
que el
>> mensaje
>> >> >> >>>> contenga. The information contained in this email is
>> >> addressed to its
>> >> >> >>>> recipient only and may contain confidential
information,
>> >> privileged
>> >> >> >>>> material or information protected by copyright. Its
>> >> prohibited any
>> >> >> >>>> copy, use, improper retention, modification,
dissemination,
>> >> >> >>>> distribution or total or partial reproduction. If you
>> receive
>> >> this
>> >> >> >>>> message by error, please contact the sender and delete
>> it. The
>> >> >> >>>> information contained herein is the sole
responsibility of
>> >> the sender
>> >> >> >>>> therefore Universidad EAFIT is not responsible for
what the
>> >> message
>> >> >> >>>> contains.
>> >> >> >>> _______________________________________________
>> >> >> >>> LLVM Developers mailing list
>> >> >> >>> llvm-dev@lists.llvm.org
>> >> >> >>>
>> >> >>
>> >>
>>
llvm-dev Info Page
>> >> >> >>>
>> >> >> >>> La información contenida en este correo electrónico está
>> dirigida
>> >> >> >>> únicamente a su destinatario y puede contener
información
>> >> >> >>> confidencial, material privilegiado o información
>> protegida por
>> >> >> >>> derecho de autor. Está prohibida cualquier copia,
>> utilización,
>> >> >> >>> indebida retención, modificación, difusión,
distribución o
>> >> >> >>> reproducción total o parcial. Si usted recibe este
>> mensaje por
>> >> error,
>> >> >> >>> por favor contacte al remitente y elimínelo. La
>> información aquí
>> >> >> >>> contenida es responsabilidad exclusiva de su
remitente por lo
>> >> tanto la
>> >> >> >>> Universidad EAFIT no se hace responsable de lo que
el mensaje
>> >> >> >>> contenga. The information contained in this email is
>> addressed
>> >> to its
>> >> >> >>> recipient only and may contain confidential information,
>> >> privileged
>> >> >> >>> material or information protected by copyright. Its
>> prohibited any
>> >> >> >>> copy, use, improper retention, modification,
dissemination,
>> >> >> >>> distribution or total or partial reproduction. If you
>> receive this
>> >> >> >>> message by error, please contact the sender and delete
>> it. The
>> >> >> >>> information contained herein is the sole responsibility
>> of the
>> >> sender
>> >> >> >>> therefore Universidad EAFIT is not responsible for
what the
>> >> message
>> >> >> >>> contains.
>> >> >> >> La información contenida en este correo electrónico está
>> dirigida
>> >> >> únicamente a su destinatario y puede contener información
>> confidencial,
>> >> >> material privilegiado o información protegida por derecho de
>> autor. Está
>> >> >> prohibida cualquier copia, utilización, indebida retención,
>> >> >> modificación, difusión, distribución o reproducción total o
>> parcial. Si
>> >> >> usted recibe este mensaje por error, por favor contacte al
>> remitente y
>> >> >> elimínelo. La información aquí contenida es responsabilidad
>> exclusiva de
>> >> >> su remitente por lo tanto la Universidad EAFIT no se hace
>> responsable de
>> >> >> lo que el mensaje contenga. The information contained in this
>> email is
>> >> >> addressed to its recipient only and may contain confidential
>> >> >> information, privileged material or information protected by
>> copyright.
>> >> >> Its prohibited any copy, use, improper retention,
modification,
>> >> >> dissemination, distribution or total or partial reproduction.
>> If you
>> >> >> receive this message by error, please contact the sender and
>> delete it.
>> >> >> The information contained herein is the sole
responsibility of the
>> >> >> sender therefore Universidad EAFIT is not responsible for
what the
>> >> >> message contains.
>> >> >> >
>> >> >>
>> >> >>
>> >> >> La información contenida en este correo electrónico está
dirigida
>> >> únicamente a su destinatario y puede contener información
confidencial,
>> >> material privilegiado o información protegida por derecho de
autor. Está
>> >> prohibida cualquier copia, utilización, indebida retención,
>> >> modificación, difusión, distribución o reproducción total o
parcial. Si
>> >> usted recibe este mensaje por error, por favor contacte al
remitente y
>> >> elimínelo. La información aquí contenida es responsabilidad
exclusiva de
>> >> su remitente por lo tanto la Universidad EAFIT no se hace
responsable de
>> >> lo que el mensaje contenga. The information contained in this
email is
>> >> addressed to its recipient only and may contain confidential
>> >> information, privileged material or information protected by
copyright.
>> >> Its prohibited any copy, use, improper retention, modification,
>> >> dissemination, distribution or total or partial reproduction.
If you
>> >> receive this message by error, please contact the sender and
delete it.
>> >> The information contained herein is the sole responsibility of the
>> >> sender therefore Universidad EAFIT is not responsible for what the
>> >> message contains.
>> >> >
>> >>
>> >>
>> >> La información contenida en este correo electrónico está dirigida
>> únicamente a su destinatario y puede contener información confidencial,
>> material privilegiado o información protegida por derecho de autor. Está
>> prohibida cualquier copia, utilización, indebida retención,
>> modificación, difusión, distribución o reproducción total o parcial. Si
>> usted recibe este mensaje por error, por favor contacte al remitente y
>> elimínelo. La información aquí contenida es responsabilidad exclusiva de
>> su remitente por lo tanto la Universidad EAFIT no se hace responsable de
>> lo que el mensaje contenga. The information contained in this email is
>> addressed to its recipient only and may contain confidential
>> information, privileged material or information protected by copyright.
>> Its prohibited any copy, use, improper retention, modification,
>> dissemination, distribution or total or partial reproduction. If you
>> receive this message by error, please contact the sender and delete it.
>> The information contained herein is the sole responsibility of the
>> sender therefore Universidad EAFIT is not responsible for what the
>> message contains.
>> >
>>
>>
>> La información contenida en este correo electrónico está dirigida
únicamente a su destinatario y puede contener información confidencial,
material privilegiado o información protegida por derecho de autor. Está
prohibida cualquier copia, utilización, indebida retención,
modificación, difusión, distribución o reproducción total o parcial. Si
usted recibe este mensaje por error, por favor contacte al remitente y
elimínelo. La información aquí contenida es responsabilidad exclusiva de
su remitente por lo tanto la Universidad EAFIT no se hace responsable de
lo que el mensaje contenga. The information contained in this email is
addressed to its recipient only and may contain confidential
information, privileged material or information protected by copyright.
Its prohibited any copy, use, improper retention, modification,
dissemination, distribution or total or partial reproduction. If you
receive this message by error, please contact the sender and delete it.
The information contained herein is the sole responsibility of the
sender therefore Universidad EAFIT is not responsible for what the
message contains.

Hi Johannes,

In order to create the testcases with the FIXMEs showing how it should look after the transformation, I've been trying to come up with a way of splitting the runtime call __tgt_target_teams based on Silei's recent patch D77005. I like the idea of using its asynchronous version __tgt_target_teams_nowait, but I think we would have to modify it, either to return the async_info struct:

__tgt_async_info __tgt_target_teams_nowait(...) // This would break the current canonical way of returning a status code.

or to receive it as a pointer and modify its content, then waiting on that async_info object:

int __tgt_target_teams_nowait(..., __tgt_async_info *async_info = nullptr)

The first way leads to modify the Clang's code generation, and I'm not sure if that is a good idea. Concerning the second way, I'm not sure if having a default argument would lead also to modify the Clang's code generation.

Another possible solution would be to implement another runtime function, avoiding the need to modify the _nowait version. This would lead to implement new runtime functions for each runtime call that implies data transfers, which is perhaps worse. I don't know which approach is better, should I look for a way of modifying Clang's code generation in case of needed?. I'm trying to be as less disruptive as possible. Can you give me some advice?.

Thank you.

Hamilton.

Hi Johannes,

>
> In order to create the testcases with the FIXMEs showing how it
should look after the transformation, I've been trying to come up with a
way of splitting the runtime call __tgt_target_teams based on Silei's
recent patch D77005. I like the idea of using its asynchronous version
__tgt_target_teams_nowait, but I think we would have to modify it,
either to return the async_info struct:
>
> __tgt_async_info __tgt_target_teams_nowait(...) // This would break
the current canonical way of returning a status code.
>
> or to receive it as a pointer and modify its content, then waiting on
that async_info object:
>
> int __tgt_target_teams_nowait(..., __tgt_async_info *async_info =
nullptr)
>
> The first way leads to modify the Clang's code generation, and I'm
not sure if that is a good idea. Concerning the second way, I'm not sure
if having a default argument would lead also to modify the Clang's code
generation.
>
> Another possible solution would be to implement another runtime
function, avoiding the need to modify the _nowait version. This would
lead to implement new runtime functions for each runtime call that
implies data transfers, which is perhaps worse. I don't know which
approach is better, should I look for a way of modifying Clang's code
generation in case of needed?. I'm trying to be as less disruptive as
possible. Can you give me some advice?.

We are in the process of moving clangs OpenMP code generation into the
OpenMPIRBuilder. While doing so we refactor/replace/rewrite stuff as we
see fit. Thus, changing Clang, or better writing a different code
generation in the OpenMPIRBuilder, is always an option.

That said, I was expecting to create two new runtime calls per existing
one but I have not really thought this through. What I imagined is
something vaguely looking like this:

Current:
int omp_transfer_data(device, src, dst, size);

New:

// In the runtime {
handle omp_transfer_data_issue(device, src, dst, size) {
   // some logic that was in omp_transfer_data before creating and using
   // handle
   return some_handle; // probably extended __tgt_async_info, can be an
                       // argument as well.
}
int omp_transfer_data_wait(device, src, dst, size, handle) {
   // some logic that was in omp_transfer_data before using handle
   return error_code;
}
// }

// In the runtime or in the header
int omp_transfer_data(device, src, dst, size) {
   handle = omp_transfer_data_issue(...);
   return omp_transfer_data_wait(..., handle);
}

Now we can inline omp_transfer_data or replace it by the two calls to
make the issue and wait explicit.

Long story short, we will most certainly need new functions, how they
look is up to you initially. Look at the other functions for inspiration
and if something is off we'll let you know during code review :wink:

Does that help? Feel free to ask more questions!

Cheers,
Johannes

P.S. Tomorrow there is an OpenMP in LLVM bi-weekly call which might be
interesting (though it might be too early):
http://lists.llvm.org/pipermail/openmp-dev/2020-April/003292.html

> Thank you.
>
> Hamilton.
>
>> > Hi Johannes,
>> >
>> > I've been thinking on how to implement a solution for the
problem and
>> > I have some questions:
>> >
>> > 1. You mention in the problem description that the current
>> > asynchronous functions can be used, and probably modified. The
problem
>> > is that those runtime functions with _nowait at the end are not
>> > asynchronous. Looking at their implementations, they just wait for
>> > other tasks is there are pending tasks, and then call its
synchronous
>> > versions. So, is it a good idea to wrap them in functions that call
>> > them as asynchronous (using std::async), then return an object
to wait
>> > on?. Then, waiting on the object returned.
>>
>> Take a look at D77005. Shilei introduced some real askync versions for a
>> different reason already. You can build on top of that patch and extend
>> it if needed further.
>>
>>
>> > 2. I'm thinking on replacing the data transfer inside
>> > __tgt_target_teams with __tgt_target_data_begin. The problem
with this
>> > is that __tgt_target_data_begin does not initialize [first]private
>> > variables, and doesn't check for lambda mappings. Therefore, my idea
>> > to go around this problem is to create another wrapper function F,
>> > that calls __tgt_target_data_begin, and executes the necessary code
>> > for lambda mappings check and [first]private variables, and finally,
>> > calling F as the asynchronous function. What is good about this, is
>> > that the code concerning data transfer inside __tgt_target_teams can
>> > be replaced with F.
>>
>> Interesting. I'll need to see small examples. For now, focus on the
>> simple case, we'll get to the rest later. However, we should really
>> start by creating test cases that show the initial code and FIXMEs that
>> describe how we want it to look like.
>>
>>
>> > 3. Also, I was wondering if I must to be aware of (split them into
>> > issue and wait) all the other constructs that imply data transfers
>> > (update, data enter/exit), or just focus on
__tgt_target[_nowait] and
>> > __tgt_target_teams[_nowait].
>>
>> We want to eventually split all transfers into two but let's start with
>> a single one.
>>
>>
>> > 4. This might be a dumb question, but just to make sure. Do I
have to
>> > move the wait call into the offloaded region just before the data is
>> > used (because I think this would imply analyzing the offloaded
>> > bitcode, which is in another module)?, or move it just before the
>> > kernel is executed?.
>>
>> We split, then we move them apart. The cannot be moved over instructions
>> that might change behavior if we move them, e.g. an aliasing memory
>> operation or a kernel call that uses the data.
>>
>>
>> > 5. Do you have a template for the GSoC proposal. That is, is it
plain
>> > text or pdf or ...?. Also, do you have sort of a structure that we
>> > have to follow?, for example, Abstract -> Benefits -> Deliverables
>> > ..., or do we have freedom on that?. Your post in hte GSoC page only
>> > mentions that we must include our background, but nothing about the
>> > format of the proposal.
>>
>> I do not have one. I thought they might or they have an FAQ that
>> describes what to do. I know they have some required things, so check
>> their webpage. I think you have to use google docs, I might be wrong.
>> Your categories look good, make the deliverables not too specific but
>> describe more ideas to work on. We'll make tasks as we go as you are not
>> the only one working on this general area.
>>
>> Feel free to send me a draft of the proposal for feedback.
>>
>> Cheers,
>> Johannes
>>
>> > Thank you.
>> >
>> > Hamilton.
>> >
>> >
>> > On``` 25/03/20 11:41 a. m., Doerfert, Johannes wrote:
>> >> > Hi Johannes,
>> >> >
>> >> > Sorry about the failing tests. I completely forgot about
running all
>> >> > of them, I only ran the tests for the openmp ipo transforms.
>> >>
>> >> No worries, it happens.
>> >>
>> >>
>> >> > I found the GH issue in cache. I can post it here if you want.
>> >>
>> >> Yes, please.
>> >>
>> >>
>> >> > I analyzed these runtime calls,
__tgt_target_teams_[nowait]. What I
>> >> > don't understand about the problem we want to solve
(splitting the
>> >> > memory transfer into issue and wait and legally moving code
>> inside) is
>> >> > that the memory transfer is made inside those runtime calls.
>> >> > Therefore, the runtime calls to the memory transfer cannot be
>> analyzed
>> >> > in the IR. Could you please explain a bit further if the memory
>> >> > transfer should split in the IR? or split
>> __tgt_target_teams_[nowait]
>> >> > in such a way that the memory transfer is made in the first
>> part, and
>> >> > the wait in the second part.
>> >>
>> >> Exactly what you said in the last sentence. Split the runtime
call into
>> >> two, issue and wait. Then we move issue "up" the execution path
and wait
>> >> "down".
>> >>
>> >> Cheers,
>> >> Johannes
>> >>
>> >>
>> >> > Thanks.
>> >> >
>> >> > Hamilton.
>> >> >
>> >> >>
>> >> >> > Hi Johannes,
>> >> >> >
>> >> >> > Yes, I'm highly interested, I already started. It
would be nice
>> >> if you
>> >> >> > post the details again. I didn't take note of them, I was
>> always
>> >> >> > looking at the issue in GH.
>> >> >>
>> >> >> I asked if we can make the readonly but unsure yet. I can
write
>> it again
>> >> >> but it'll take ma some time.
>> >> >>
>> >> >>
>> >> >> > The problem I'm facing right now is this:
>> >> >> > The first task is identifying OpenMP API calls referred to
>> >> device memory
>> >> >> > allocation. I first understood how the compiler driver
>> >> orchestrates the
>> >> >> > process of offloading OpenMP regions to devices, and
now I'm
>> >> looking at
>> >> >> > the step of the pipeline that generates LLVM IR for both
>> host and
>> >> >> > device. I wrote a simple C program that uses OpenMP to
map a
>> >> local array
>> >> >> > to the device main memory. Then I'm looking at the
device and
>> >> host LLVM
>> >> >> > IR representation of this program, but I can't find any
>> runtime call
>> >> >> > referred to device memory allocation. In the host IR I see
>> that the
>> >> >> > array is passed to the outlined function and stored in a
>> >> register inside
>> >> >> > the outlined function, but there's no memory
allocation runtime
>> >> call.
>> >> >> > And in the device IR the array is passed to the main
offloading
>> >> >> > function.
>> >> >> > Any suggestion or documentation on how that memory
transfer is
>> >> >> > achieved?, and how to track its API calls. It might be
that
>> I'm not
>> >> >> > using the proper omp directive, so here is C code:
>> >> >>
>> >> >> Take a look at
>> >>
>>

>> >> , the mapping request,
>> >> >>
>> >> >> which may or may not cause data-transfers, are part of the
call to
>> >> >> __tgt_target_teams[_nowait] and
>> __tgt_target_data_update[_nowait]. Start
>> >> >> by reading the impl. of those two functions in the openmp
runtime
>> >> >> (llvm-project/openmp/...).
>> >> >>
>> >> >> Let me know if this helps and once you have new questions.
>> >> >>
>> >> >> Cheers,
>> >> >> Johannes
>> >> >>
>> >> >> P.S. I merged your first patch. You didn't run all test
before.
>> I did
>> >> >> and noticed that we failed one because of type mismatches.
>> Fixed them
>> >> >> before I merged your patch.
>> >> >>
>> >> >> >
>> >> >> > #include <omp.h>
>> >> >> > #include <stdio.h>
>> >> >> >
>> >> >> > #define size 1000000
>> >> >> > double a[size];
>> >> >> >
>> >> >> > int main() {
>> >> >> >
>> >> >> > #pragma omp target map(to: a[0:size])
>> >> >> > #pragma omp teams distribute
>> >> >> > for (int i = 0; i < size; ++i) {
>> >> >> > a[i] = a[i]*(i*i)/2;
>> >> >> > }
>> >> >> >
>> >> >> > return 0;
>> >> >> > }
>> >> >> >
>> >> >> >
>> >> >> > PD: The command to compile is this:
>> >> >> > clang -v -save-temps \
>> >> >> > -I/path/to/llvm/release/9.x/include -fopenmp \
>> >> >> > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
>> >> >> > -march=sm_61 memory_transfer.c -o memory_transfer
>> >> >> >
>> >> >> > Thanks.
>> >> >> > Hamilton.
>> >> >> >
>> >> >> >> > Ahhh I understand. So I believe the project you
posted
>> is still
>> >> >> >> opened. Do you any advice on the question I posted
before?.
>> >> >> >>
>> >> >> >> Yes. Still open. You should use this as the main part of
>> the GSoC
>> >> >> >> proposal (if you are interested). Can you repeat the
question?
>> >> >> >>
>> >> >> >> >
>> >> >> >> > Thank you.
>> >> >> >> >
>> >> >> >> >> I did not. The issues sections on GH was removed.
>> >> >> >> >>
>> >> >> >> >> ________________________________________
>> >> >> >> >> From: Hamilton Tobon Mosquera <htobonm@eafit.edu.co>
>> >> >> >> >> Sent: Sunday, March 22, 2020 12:21
>> >> >> >> >> To: llvm-dev@lists.llvm.org
>> >> >> >> >> Cc: Doerfert, Johannes
>> >> >> >> >> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >> >> >>
>> >> >> >> >> Hi there Johannes!,
>> >> >> >> >>
>> >> >> >> >> I was checking out the issue (about latency
hiding when
>> >> >> transferring
>> >> >> >> >> memory from host to device) you posted about a week
>> ago, but
>> >> >> it's not
>> >> >> >> >> there anymore. Did you remove it?.
>> >> >> >> >>
>> >> >> >> >> Thanks.
>> >> >> >> >>
>> >> >> >> >> Hamilton.
>> >> >> >> >>
>> >> >> >> >>> Hi there!,
>> >> >> >> >>>
>> >> >> >> >>> Right now I'm starting with the project proposed by
>> >> Johannes here:
>> >> >> >> >>>
>> >> >> >>
>> >> >>
>> >>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub.
>> >> >> >> >>>
>> >> >> >> >>>
>> >> >> >> >>> The first task is identifying OpenMP API calls
referred
>> >> to device
>> >> >> >> memory
>> >> >> >> >>> allocation. I first understood how the compiler
driver
>> >> >> orchestrates the
>> >> >> >> >>> process of offloading OpenMP regions to
devices, and
>> now I'm
>> >> >> looking at
>> >> >> >> >>> the step of the pipeline that generates LLVM IR
for both
>> >> host and
>> >> >> >> >>> device. I wrote a simple C program that uses
OpenMP to
>> >> map a local
>> >> >> >> array
>> >> >> >> >>> to the device main memory. Then I'm looking at the
>> device and
>> >> >> host LLVM
>> >> >> >> >>> IR representation of this program, but I can't
find any
>> >> >> runtime call
>> >> >> >> >>> referred to device memory allocation. In the
host IR
>> I see
>> >> >> that the
>> >> >> >> >>> array is passed to the outlined function and
stored in a
>> >> register
>> >> >> >> inside
>> >> >> >> >>> the outlined function, but there's no memory
allocation
>> >> >> runtime call.
>> >> >> >> >>> And in the device IR the array is passed to the
main
>> >> offloading
>> >> >> >> >>> function. Any suggestion or documentation on how
>> that memory
>> >> >> >> transfer is
>> >> >> >> >>> achieved?, and how to track its API calls. It might
>> be that
>> >> >> I'm not
>> >> >> >> >>> using the proper omp directive, so here is C code:
>> >> >> >> >>>
>> >> >> >> >>>
>> >> >> >> >>> #include <omp.h>
>> >> >> >> >>> #include <stdio.h>
>> >> >> >> >>>
>> >> >> >> >>> #define size 1000000
>> >> >> >> >>> double a[size];
>> >> >> >> >>>
>> >> >> >> >>> int main() {
>> >> >> >> >>>
>> >> >> >> >>> #pragma omp target map(to: a[0:size])
>> >> >> >> >>> #pragma omp teams distribute
>> >> >> >> >>> for (int i = 0; i < size; ++i) {
>> >> >> >> >>> a[i] = a[i]*(i*i)/2;
>> >> >> >> >>> }
>> >> >> >> >>> return 0;
>> >> >> >> >>>
>> >> >> >> >>> PD: The command to compile is this: clang -v
-save-temps
>> >> >> >> >>> -I/path/to/llvm/release/9.x/include -fopenmp
>> >> >> >> >>> -fopenmp-targets=nvptx64-nvidia-cuda
-Xopenmp-target
>> >> -march=sm_61
>> >> >> >> >>> memory_transfer.c -o memory_transfer
>> >> >> >> >>>
>> >> >> >> >>> Thank you all.
>> >> >> >> >>>
>> >> >> >> >>> Hamilton.
>> >> >> >> >>>
>> >> >> >> >>>
>> >> >> >> >>>> Hi Hamilton,
>> >> >> >> >>>>
>> >> >> >> >>>> I personally don't believe the effort to make
the IR
>> >> >> >> >>>> "parallelism-aware" is worth it right now.
>> >> >> >> >>>> What I propose is something else, see for example
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fyoutu.be%2FzfiHaPaoQPc&amp;data=01|01|htobonm%40eafit.edu.co|3e9ac7a8643c4dbe405e08d7d38dfd75|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=hN9WGEO%2F3quio7A9tvtIDHxznURuITYVBszFSmqM5M8%3D&amp;reserved=0
>> >> >> >> >>>> and the OpenMPOpt pass.
>> >> >> >> >>>> There are multiple bigger tasks towards better
>> >> offloading, one of
>> >> >> >> >>>> which is described here
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
[OpenMP] Memory transfer optimizations · Issue #180 · llvm/llvm-project · GitHub
>> >> >> >> >>>> It might be interesting to add the transfer
>> functions used
>> >> >> for memory
>> >> >> >> >>>> transfer to the OMPKinds.def file so they are
known
>> in the
>> >> >> >> >>>> OpenMPOpt.cpp pass.
>> >> >> >> >>>> Another small task would be to eliminate trivially
>> redundant
>> >> >> ones,
>> >> >> >> >>>> that is if you have copy X from A to B directly
>> followed by
>> >> >> copy X
>> >> >> >> >>>> from B to A, remove both.
>> >> >> >> >>>>
>> >> >> >> >>>> Does this make sense?
>> >> >> >> >>>>
>> >> >> >> >>>> Cheers,
>> >> >> >> >>>> Johannes
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> ---------------------------------------
>> >> >> >> >>>> Johannes Doerfert
>> >> >> >> >>>> Researcher
>> >> >> >> >>>>
>> >> >> >> >>>> Argonne National Laboratory
>> >> >> >> >>>> Lemont, IL 60439, USA
>> >> >> >> >>>>
>> >> >> >> >>>> jdoerfert@anl.gov
>> >> >> >> >>>>
>> >> >> >> >>>> ________________________________________
>> >> >> >> >>>> From: Hamilton Tobon Mosquera
<htobonm@eafit.edu.co>
>> >> >> >> >>>> Sent: Wednesday, March 18, 2020 19:14
>> >> >> >> >>>> To: llvm-dev@lists.llvm.org
>> >> >> >> >>>> Cc: Doerfert, Johannes
>> >> >> >> >>>> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >> >> >>>>
>> >> >> >> >>>> Hi there!,
>> >> >> >> >>>>
>> >> >> >> >>>> I think I'm done with this patch: D76058. If
that is
>> >> everything
>> >> >> >> >>>> concerning that, could you please assign me
another
>> task.
>> >> >> >> >>>>
>> >> >> >> >>>> I've been reading about Concurrent Control
Flow Graphs,
>> >> which
>> >> >> would
>> >> >> >> >>>> be awesome for LLVM because it would be kind of a
>> starting
>> >> >> point for
>> >> >> >> >>>> optimizations for parallel programs in
general. The
>> thing is
>> >> >> that it
>> >> >> >> >>>> seems more complicated than it is presented in the
>> papers, I
>> >> >> think
>> >> >> >> >>>> that's why it's not implemented in LLVM yet. So, I
>> think
>> >> I'll try
>> >> >> >> >>>> searching for something else. Nevertheless,
don't you
>> >> think that
>> >> >> >> >>>> implementing this is worth the effort?. The
papers I'm
>> >> >> reading are
>> >> >> >> >>>> from 1990s, I don't understand why they aren't
widely
>> >> >> implemented in
>> >> >> >> >>>> big compilers xd.
>> >> >> >> >>>>
>> >> >> >> >>>> I saw this commit from you:
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
⚙ D29250 [PIR] Add fork, join and halt instructions.
>> >> >> >> >>>> Is there something I can help with concerning
this?.
>> >> >> >> >>>>
>> >> >> >> >>>> I would like to do contribute to the OpenMP target
>> >> offloading. Is
>> >> >> >> >>>> there anything I can help with?.
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> Hi there!,
>> >> >> >> >>>>
>> >> >> >> >>>> As part of my application process to the next
GSoC I'm
>> >> >> working on the
>> >> >> >> >>>> TODO in OpenMPOpt.cpp line 437:
>> >> >> >> >>>>
>> >> >> >> >>>> // TODO: We should validate the declaration
>> agains the
>> >> >> types we
>> >> >> >> >>>> expect.
>> >> >> >> >>>>
>> >> >> >> >>>> Link:
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fblob%2Fmaster%2Fllvm%2Flib%2FTransforms%2FIPO%2FOpenMPOpt.cpp%23L437&amp;data=01|01|htobonm%40eafit.edu.co|3e9ac7a8643c4dbe405e08d7d38dfd75|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=gAMSEONp%2F8cMDuO0IckxHPXbBnUllJMWVfLHwa6UsBc%3D&amp;reserved=0
>> >> >> >> >>>>
>> >> >> >> >>>> I have a question. When there is a mismatch in the
>> types
>> >> >> (return type
>> >> >> >> >>>> or argument types) between the function
declaration
>> >> found and the
>> >> >> >> >>>> runtime library function, what should I do?.
Show an
>> >> error or a
>> >> >> >> >>>> warning?, continuing with the pass or exiting the
>> program?.
>> >> >> >> >>>>
>> >> >> >> >>>> Thanks.
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> Hi Hamilton,
>> >> >> >> >>>>
>> >> >> >> >>>> Thanks for your interest in LLVM!
>> >> >> >> >>>>
>> >> >> >> >>>> IMHO you have a good start. I'll try to help
by CC'ing
>> >> >> Johannes. Note
>> >> >> >> >>>> that it's difficult for LLVM developers
>> >> >> >> >>>> (including GSoC mentors and including Johannes) to
>> keep up
>> >> >> with every
>> >> >> >> >>>> discussion on llvm-dev.
>> >> >> >> >>>> So, you can CC them to make their life easier. :slight_smile:
>> >> >> >> >>>>
>> >> >> >> >>>> Johannes, I see a lot of people interested in this
>> project.
>> >> >> Maybe a
>> >> >> >> >>>> list of starting points, patches etc.
>> >> >> >> >>>> could help.
>> >> >> >> >>>>
>> >> >> >> >>>> Best,
>> >> >> >> >>>> Stefanos
>> >> >> >> >>>>
>> >> >> >> >>>> Στις Τρί, 10 Μαρ 2020 στις 1:04 π.μ., ο/η Hamilton
>> Tobon
>> >> >> Mosquera via
>> >> >> >> >>>> llvm-dev
>> >> >> <llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>>
>> >> >> >> >>>> έγραψε:
>> >> >> >> >>>> Greetings,
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> I'm an LLVM newcomer interested in
participating in the
>> >> next GSoC
>> >> >> >> under
>> >> >> >> >>>> the project "Improve parallelism-aware
analyses and
>> >> >> >> optimizations". I've
>> >> >> >> >>>> already read "Compiler Optimizations for
OpenMP" and
>> >> "Compiler
>> >> >> >> >>>> Optimizations for Parallel Programs" both by
Doerfert
>> >> and Finkel.
>> >> >> >> Also,
>> >> >> >> >>>> I've watched their two LLVM meeting developers
>> conferences
>> >> >> >> "Representing
>> >> >> >> >>>> parallelism within LLVM" and "Optimizing
>> indirections, using
>> >> >> >> >>>> abstractions without ...", so I have some
>> background in the
>> >> >> problems
>> >> >> >> >>>> they are trying to tackle. Also, I'm close to
solve
>> one of
>> >> >> the bugs
>> >> >> >> >>>> posted in bugzilla, which has given me a broader
>> idea of
>> >> how the
>> >> >> >> OpenMP
>> >> >> >> >>>> pragmas are translated. The bus is this:
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&amp;data=01|01|htobonm%40eafit.edu.co|3e9ac7a8643c4dbe405e08d7d38dfd75|99f7b55e9cbe467b8143919782918afb|0&amp;sdata=Wye5YuFLVUPPESqTRnwsZtimjF0wdUxDetwoCNCib64%3D&amp;reserved=0.
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> I tried contacting you Johannes but with your
>> Argonne email,
>> >> >> I suppose
>> >> >> >> >>>> that's why you couldn't answer. Could you tell
me more
>> >> about the
>> >> >> >> >>>> project?, what variants do you have?, or what
can I
>> focus my
>> >> >> >> application
>> >> >> >> >>>> on?. Should I start with what you recommended to
>> Emanuel?.
>> >> >> >> >>>>
>> >> >> >> >>>> Thank you.
>> >> >> >> >>>>
>> >> >> >> >>>> _______________________________________________
>> >> >> >> >>>> LLVM Developers mailing list
>> >> >> >> >>>>
llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
llvm-dev Info Page
>> >> >> >> >>>>
>> >> >> >> >>>> La información contenida en este correo
electrónico
>> está
>> >> dirigida
>> >> >> >> >>>> únicamente a su destinatario y puede contener
>> información
>> >> >> >> >>>> confidencial, material privilegiado o información
>> >> protegida por
>> >> >> >> >>>> derecho de autor. Está prohibida cualquier copia,
>> >> utilización,
>> >> >> >> >>>> indebida retención, modificación, difusión,
>> distribución o
>> >> >> >> >>>> reproducción total o parcial. Si usted recibe este
>> mensaje
>> >> >> por error,
>> >> >> >> >>>> por favor contacte al remitente y elimínelo. La
>> >> información aquí
>> >> >> >> >>>> contenida es responsabilidad exclusiva de su
remitente
>> >> por lo
>> >> >> tanto
>> >> >> >> >>>> la Universidad EAFIT no se hace responsable de lo
>> que el
>> >> mensaje
>> >> >> >> >>>> contenga. The information contained in this
email is
>> >> >> addressed to its
>> >> >> >> >>>> recipient only and may contain confidential
>> information,
>> >> >> privileged
>> >> >> >> >>>> material or information protected by
copyright. Its
>> >> >> prohibited any
>> >> >> >> >>>> copy, use, improper retention, modification,
>> dissemination,
>> >> >> >> >>>> distribution or total or partial reproduction.
If you
>> >> receive
>> >> >> this
>> >> >>�� >> >>>> message by error, please contact the sender
and delete
>> >> it. The
>> >> >> >> >>>> information contained herein is the sole
>> responsibility of
>> >> >> the sender
>> >> >> >> >>>> therefore Universidad EAFIT is not responsible for
>> what the
>> >> >> message
>> >> >> >> >>>> contains.
>> >> >> >> >>>> La información contenida en este correo
electrónico
>> está
>> >> dirigida
>> >> >> >> >>>> únicamente a su destinatario y puede contener
>> información
>> >> >> >> >>>> confidencial, material privilegiado o información
>> >> protegida por
>> >> >> >> >>>> derecho de autor. Está prohibida cualquier copia,
>> >> utilización,
>> >> >> >> >>>> indebida retención, modificación, difusión,
>> distribución o
>> >> >> >> >>>> reproducción total o parcial. Si usted recibe este
>> mensaje
>> >> >> por error,
>> >> >> >> >>>> por favor contacte al remitente y elimínelo. La
>> >> información aquí
>> >> >> >> >>>> contenida es responsabilidad exclusiva de su
remitente
>> >> por lo
>> >> >> tanto
>> >> >> >> >>>> la Universidad EAFIT no se hace responsable de lo
>> que el
>> >> mensaje
>> >> >> >> >>>> contenga. The information contained in this
email is
>> >> >> addressed to its
>> >> >> >> >>>> recipient only and may contain confidential
>> information,
>> >> >> privileged
>> >> >> >> >>>> material or information protected by
copyright. Its
>> >> >> prohibited any
>> >> >> >> >>>> copy, use, improper retention, modification,
>> dissemination,
>> >> >> >> >>>> distribution or total or partial reproduction.
If you
>> >> receive
>> >> >> this
>> >> >> >> >>>> message by error, please contact the sender
and delete
>> >> it. The
>> >> >> >> >>>> information contained herein is the sole
>> responsibility of
>> >> >> the sender
>> >> >> >> >>>> therefore Universidad EAFIT is not responsible for
>> what the
>> >> >> message
>> >> >> >> >>>> contains.
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> _______________________________________________
>> >> >> >> >>>> LLVM Developers mailing list
>> >> >> >> >>>>
llvm-dev@lists.llvm.org<mailto:llvm-dev@lists.llvm.org>
>> >> >> >> >>>>
>> >> >> >>
>> >> >>
>> >>
>>
llvm-dev Info Page
>> >> >> >> >>>>
>> >> >> >> >>>>
>> >> >> >> >>>> La información contenida en este correo
electrónico
>> está
>> >> dirigida
>> >> >> >> >>>> únicamente a su destinatario y puede contener
>> información
>> >> >> >> >>>> confidencial, material privilegiado o información
>> >> protegida por
>> >> >> >> >>>> derecho de autor. Está prohibida cualquier copia,
>> >> utilización,
>> >> >> >> >>>> indebida retención, modificación, difusión,
>> distribución o
>> >> >> >> >>>> reproducción total o parcial. Si usted recibe este
>> mensaje
>> >> >> por error,
>> >> >> >> >>>> por favor contacte al remitente y elimínelo. La
>> >> información aquí
>> >> >> >> >>>> contenida es responsabilidad exclusiva de su
remitente
>> >> por lo
>> >> >> tanto
>> >> >> >> >>>> la Universidad EAFIT no se hace responsable de lo
>> que el
>> >> mensaje
>> >> >> >> >>>> contenga. The information contained in this
email is
>> >> >> addressed to its
>> >> >> >> >>>> recipient only and may contain confidential
>> information,
>> >> >> privileged
>> >> >> >> >>>> material or information protected by
copyright. Its
>> >> >> prohibited any
>> >> >> >> >>>> copy, use, improper retention, modification,
>> dissemination,
>> >> >> >> >>>> distribution or total or partial reproduction.
If you
>> >> receive
>> >> >> this
>> >> >> >> >>>> message by error, please contact the sender
and delete
>> >> it. The
>> >> >> >> >>>> information contained herein is the sole
>> responsibility of
>> >> >> the sender
>> >> >> >> >>>> therefore Universidad EAFIT is not responsible for
>> what the
>> >> >> message
>> >> >> >> >>>> contains.
>> >> >> >> >>> _______________________________________________
>> >> >> >> >>> LLVM Developers mailing list
>> >> >> >> >>> llvm-dev@lists.llvm.org
>> >> >> >> >>>
>> >> >> >>
>> >> >>
>> >>
>>
llvm-dev Info Page
>> >> >> >> >>>
>> >> >> >> >>> La información contenida en este correo
electrónico está
>> >> dirigida
>> >> >> >> >>> únicamente a su destinatario y puede contener
>> información
>> >> >> >> >>> confidencial, material privilegiado o información
>> >> protegida por
>> >> >> >> >>> derecho de autor. Está prohibida cualquier copia,
>> >> utilización,
>> >> >> >> >>> indebida retención, modificación, difusión,
>> distribución o
>> >> >> >> >>> reproducción total o parcial. Si usted recibe este
>> >> mensaje por
>> >> >> error,
>> >> >> >> >>> por favor contacte al remitente y elimínelo. La
>> >> información aquí
>> >> >> >> >>> contenida es responsabilidad exclusiva de su
>> remitente por lo
>> >> >> tanto la
>> >> >> >> >>> Universidad EAFIT no se hace responsable de lo que
>> el mensaje
>> >> >> >> >>> contenga. The information contained in this
email is
>> >> addressed
>> >> >> to its
>> >> >> >> >>> recipient only and may contain confidential
information,
>> >> >> privileged
>> >> >> >> >>> material or information protected by copyright. Its
>> >> prohibited any
>> >> >> >> >>> copy, use, improper retention, modification,
>> dissemination,
>> >> >> >> >>> distribution or total or partial reproduction.
If you
>> >> receive this
>> >> >> >> >>> message by error, please contact the sender and
delete
>> >> it. The
>> >> >> >> >>> information contained herein is the sole
responsibility
>> >> of the
>> >> >> sender
>> >> >> >> >>> therefore Universidad EAFIT is not responsible for
>> what the
>> >> >> message
>> >> >> >> >>> contains.
>> >> >> >> >> La información contenida en este correo
electrónico está
>> >> dirigida
>> >> >> >> únicamente a su destinatario y puede contener información
>> >> confidencial,
>> >> >> >> material privilegiado o información protegida por
derecho de
>> >> autor. Está
>> >> >> >> prohibida cualquier copia, utilización, indebida
retención,
>> >> >> >> modificación, difusión, distribución o reproducción
total o
>> >> parcial. Si
>> >> >> >> usted recibe este mensaje por error, por favor
contacte al
>> >> remitente y
>> >> >> >> elimínelo. La información aquí contenida es
responsabilidad
>> >> exclusiva de
>> >> >> >> su remitente por lo tanto la Universidad EAFIT no se hace
>> >> responsable de
>> >> >> >> lo que el mensaje contenga. The information contained
in this
>> >> email is
>> >> >> >> addressed to its recipient only and may contain
confidential
>> >> >> >> information, privileged material or information
protected by
>> >> copyright.
>> >> >> >> Its prohibited any copy, use, improper retention,
>> modification,
>> >> >> >> dissemination, distribution or total or partial
reproduction.
>> >> If you
>> >> >> >> receive this message by error, please contact the
sender and
>> >> delete it.
>> >> >> >> The information contained herein is the sole
>> responsibility of the
>> >> >> >> sender therefore Universidad EAFIT is not responsible for
>> what the
>> >> >> >> message contains.
>> >> >> >> >
>> >> >> >>
>> >> >> >>
>> >> >> >> La información contenida en este correo electrónico está
>> dirigida
>> >> >> únicamente a su destinatario y puede contener información
>> confidencial,
>> >> >> material privilegiado o información protegida por derecho de
>> autor. Está
>> >> >> prohibida cualquier copia, utilización, indebida retención,
>> >> >> modificación, difusión, distribución o reproducción total o
>> parcial. Si
>> >> >> usted recibe este mensaje por error, por favor contacte al
>> remitente y
>> >> >> elimínelo. La información aquí contenida es responsabilidad
>> exclusiva de
>> >> >> su remitente por lo tanto la Universidad EAFIT no se hace
>> responsable de
>> >> >> lo que el mensaje contenga. The information contained in this
>> email is
>> >> >> addressed to its recipient only and may contain confidential
>> >> >> information, privileged material or information protected by
>> copyright.
>> >> >> Its prohibited any copy, use, improper retention,
modification,
>> >> >> dissemination, distribution or total or partial reproduction.
>> If you
>> >> >> receive this message by error, please contact the sender and
>> delete it.
>> >> >> The information contained herein is the sole
responsibility of the
>> >> >> sender therefore Universidad EAFIT is not responsible for
what the
>> >> >> message contains.
>> >> >> >
>> >> >>
>> >> >>
>> >> >> La información contenida en este correo electrónico está
dirigida
>> >> únicamente a su destinatario y puede contener información
confidencial,
>> >> material privilegiado o información protegida por derecho de
autor. Está
>> >> prohibida cualquier copia, utilización, indebida retención,
>> >> modificación, difusión, distribución o reproducción total o
parcial. Si
>> >> usted recibe este mensaje por error, por favor contacte al
remitente y
>> >> elimínelo. La información aquí contenida es responsabilidad
exclusiva de
>> >> su remitente por lo tanto la Universidad EAFIT no se hace
responsable de
>> >> lo que el mensaje contenga. The information contained in this
email is
>> >> addressed to its recipient only and may contain confidential
>> >> information, privileged material or information protected by
copyright.
>> >> Its prohibited any copy, use, improper retention, modification,
>> >> dissemination, distribution or total or partial reproduction.
If you
>> >> receive this message by error, please contact the sender and
delete it.
>> >> The information contained herein is the sole responsibility of the
>> >> sender therefore Universidad EAFIT is not responsible for what the
>> >> message contains.
>> >> >
>> >>
>> >>
>> >> La información contenida en este correo electrónico está dirigida
>> únicamente a su destinatario y puede contener información confidencial,
>> material privilegiado o información protegida por derecho de autor. Está
>> prohibida cualquier copia, utilización, indebida retención,
>> modificación, difusión, distribución o reproducción total o parcial. Si
>> usted recibe este mensaje por error, por favor contacte al remitente y
>> elimínelo. La información aquí contenida es responsabilidad exclusiva de
>> su remitente por lo tanto la Universidad EAFIT no se hace responsable de
>> lo que el mensaje contenga. The information contained in this email is
>> addressed to its recipient only and may contain confidential
>> information, privileged material or information protected by copyright.
>> Its prohibited any copy, use, improper retention, modification,
>> dissemination, distribution or total or partial reproduction. If you
>> receive this message by error, please contact the sender and delete it.
>> The information contained herein is the sole responsibility of the
>> sender therefore Universidad EAFIT is not responsible for what the
>> message contains.
>> >
>>
>>
>> La información contenida en este correo electrónico está dirigida
únicamente a su destinatario y puede contener información confidencial,
material privilegiado o información protegida por derecho de autor. Está
prohibida cualquier copia, utilización, indebida retención,
modificación, difusión, distribución o reproducción total o parcial. Si
usted recibe este mensaje por error, por favor contacte al remitente y
elimínelo. La información aquí contenida es responsabilidad exclusiva de
su remitente por lo tanto la Universidad EAFIT no se hace responsable de
lo que el mensaje contenga. The information contained in this email is
addressed to its recipient only and may contain confidential
information, privileged material or information protected by copyright.
Its prohibited any copy, use, improper retention, modification,
dissemination, distribution or total or partial reproduction. If you
receive this message by error, please contact the sender and delete it.
The information contained herein is the sole responsibility of the
sender therefore Universidad EAFIT is not responsible for what the
message contains.

Hi Johannes,

I'm still writing the test cases with the FIXMEs that show how the runtime call should be split and moved forward and backwards. I'm starting with __tgt_target_teams and the possible situations that might occur. The patch that Shilei did allows for making the computation asynchronous too, not only the data transfer, it looks something like:

...

target_data_begin(..., __tgt_async_info a)

...

run_team_region(..., __tgt_async_info a)

target_data_end(..., __tgt_async_info a)

...

synchronize(__tgt_async_info a)

My question is, should the wait call contain target_data_run(...), that is, running the target region only when calling wait()?, or should it only call synchronize(...), that is, leaving the computation to the issue(...) function?. The first case would look something like:

issue() {

...

target_data_begin()

...

}

wait() {

run_team_region() // The second approach has this function in the issue

target_data_end() // as well as this one.

synchronize()

}

The thing is that some behaviors change given the first or second approach. For example:

heavyComputation() {

...

int a = rand();

...

int x = a % rand();

...

args[0] = &a;

issue(..., args, ...)

...

}

In that example, if the variable 'a' is shared and the 'issue' function has the 'run_team_region' call, it wouldn't be correct to move the issue past the expression 'int x = a % rand();', because a might be modified. But if the issue only transfers data, there wouldn't be any problem I guess. I think either approach has its advantages and disadvantages, for example having the computation in the issue would prevent the transfer to be moved even more forward. What do you think I should do?.

Thanks.

Hamilton.