llvm dev - Mar 2020 - GSoC Interested student

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:
> On 3/25/20 5:51 AM, Hamilton Tobon Mosquera 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.
>   >
>   > On 23/03/20 2:21 p. m., Doerfert, Johannes wrote:
>   >>
>   >> On 3/23/20 11:22 AM, Hamilton Tobon Mosquera wrote:
>   >>   > 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
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=nhtF7ZKL3iPAQEJpIgCFb0cpcPGQYRvOHuPmkxyI6H8%3D&reserved=0
> , 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.
>   >>   >
>   >>   > On 23/03/20 11:11 a. m., Doerfert, Johannes wrote:
>   >>   >> On 3/22/20 12:33 PM, Hamilton Tobon Mosquera wrote:
>   >>   >>   > 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.
>   >>   >>   >
>   >>   >>   > On 22/03/20 1:26 p. m., Doerfert, Johannes
wrote:
>   >>   >>   >> I did not. The issues sections on GH was
removed.
>   >>   >>   >>
>   >>   >>   >> ________________________________________
>   >>   >>   >> From: Hamilton Tobon Mosquera <htobonm
at eafit.edu.co>
>   >>   >>   >> Sent: Sunday, March 22, 2020 12:21
>   >>   >>   >> To: llvm-dev at 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.
>   >>   >>   >>
>   >>   >>   >> On 21/03/20 10:04 p. m., Hamilton Tobon
Mosquera via llvm-dev
>   >> wrote:
>   >>   >>   >>> Hi there!,
>   >>   >>   >>>
>   >>   >>   >>> Right now I'm starting with the
project proposed by
> Johannes here:
>   >>   >>   >>>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=wyVbkLe8jhPzEIH6zHRwZZLdPHCKXPuXQNW%2BOhMHxv4%3D&amp;reserved=0.
>   >>   >>   >>>
>   >>   >>   >>>
>   >>   >>   >>> 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.
>   >>   >>   >>>
>   >>   >>   >>>
>   >>   >>   >>> On 18/03/20 9:21 p. m., Doerfert,
Johannes wrote:
>   >>   >>   >>>> 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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&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
>   >>   >>   >>>>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=wyVbkLe8jhPzEIH6zHRwZZLdPHCKXPuXQNW%2BOhMHxv4%3D&amp;reserved=0
>   >>   >>   >>>> 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 at anl.gov
>   >>   >>   >>>>
>   >>   >>   >>>>
________________________________________
>   >>   >>   >>>> From: Hamilton Tobon Mosquera
<htobonm at eafit.edu.co>
>   >>   >>   >>>> Sent: Wednesday, March 18, 2020
19:14
>   >>   >>   >>>> To: llvm-dev at 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:
>   >>   >>   >>>>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD29250&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=%2BLHYDJp%2B%2FT1xu%2FesfcvHVUeJweV7MKGjcVE58Twmsk8%3D&amp;reserved=0.
>   >>   >>   >>>> 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?.
>   >>   >>   >>>>
>   >>   >>   >>>> On 11/03/20 8:51 a. m., Hamilton
Tobon Mosquera via llvm-dev
>   >> wrote:
>   >>   >>   >>>>
>   >>   >>   >>>> 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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=47sDHCNwfoSt%2B7HDCHg5xe7yTK2IYYPOY6k2NxvGaj8%3D&amp;reserved=0<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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&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.
>   >>   >>   >>>>
>   >>   >>   >>>>
>   >>   >>   >>>> On 10/03/20 7:12 p. m., Stefanos
Baziotis wrote:
>   >>   >>   >>>> 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. :)
>   >>   >>   >>>>
>   >>   >>   >>>> 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 at lists.llvm.org<mailto:llvm-dev at
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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=r2Gq2MnAyC8jZzMYPHkEKRnL7p2teev2K3t9%2FdtE9es%3D&amp;reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&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 at
lists.llvm.org<mailto:llvm-dev at lists.llvm.org>
>   >>   >>   >>>>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0>
>   >>   >>   >>>>
>   >>   >>   >>>> 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 at
lists.llvm.org<mailto:llvm-dev at lists.llvm.org>
>   >>   >>   >>>>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0
>   >>   >>   >>>>
>   >>   >>   >>>>
>   >>   >>   >>>> 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 at lists.llvm.org
>   >>   >>   >>>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0
>   >>   >>   >>>
>   >>   >>   >>> 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.

On 3/28/20 8:33 PM, Hamilton Tobon Mosquera 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:
 >> On 3/25/20 5:51 AM, Hamilton Tobon Mosquera 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.
 >>   >
 >>   > On 23/03/20 2:21 p. m., Doerfert, Johannes wrote:
 >>   >>
 >>   >> On 3/23/20 11:22 AM, Hamilton Tobon Mosquera wrote:
 >>   >>   > 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
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=nhtF7ZKL3iPAQEJpIgCFb0cpcPGQYRvOHuPmkxyI6H8%3D&reserved=0
 >> , 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.
 >>   >>   >
 >>   >>   > On 23/03/20 11:11 a. m., Doerfert, Johannes wrote:
 >>   >>   >> On 3/22/20 12:33 PM, Hamilton Tobon Mosquera
wrote:
 >>   >>   >>   > 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.
 >>   >>   >>   >
 >>   >>   >>   > On 22/03/20 1:26 p. m., Doerfert,
Johannes wrote:
 >>   >>   >>   >> I did not. The issues sections on GH
was removed.
 >>   >>   >>   >>
 >>   >>   >>   >>
________________________________________
 >>   >>   >>   >> From: Hamilton Tobon Mosquera
<htobonm at eafit.edu.co>
 >>   >>   >>   >> Sent: Sunday, March 22, 2020 12:21
 >>   >>   >>   >> To: llvm-dev at 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.
 >>   >>   >>   >>
 >>   >>   >>   >> On 21/03/20 10:04 p. m., Hamilton
Tobon Mosquera via
llvm-dev
 >>   >> wrote:
 >>   >>   >>   >>> Hi there!,
 >>   >>   >>   >>>
 >>   >>   >>   >>> Right now I'm starting with
the project proposed by
 >> Johannes here:
 >>   >>   >>   >>>
 >>   >>   >>
 >>   >>
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=wyVbkLe8jhPzEIH6zHRwZZLdPHCKXPuXQNW%2BOhMHxv4%3D&amp;reserved=0.
 >>   >>   >>   >>>
 >>   >>   >>   >>>
 >>   >>   >>   >>> 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.
 >>   >>   >>   >>>
 >>   >>   >>   >>>
 >>   >>   >>   >>> On 18/03/20 9:21 p. m., Doerfert,
Johannes wrote:
 >>   >>   >>   >>>> 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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&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
 >>   >>   >>   >>>>
 >>   >>   >>
 >>   >>
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=wyVbkLe8jhPzEIH6zHRwZZLdPHCKXPuXQNW%2BOhMHxv4%3D&amp;reserved=0
 >>   >>   >>   >>>> 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 at anl.gov
 >>   >>   >>   >>>>
 >>   >>   >>   >>>>
________________________________________
 >>   >>   >>   >>>> From: Hamilton Tobon Mosquera
<htobonm at eafit.edu.co>
 >>   >>   >>   >>>> Sent: Wednesday, March 18,
2020 19:14
 >>   >>   >>   >>>> To: llvm-dev at
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:
 >>   >>   >>   >>>>
 >>   >>   >>
 >>   >>
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD29250&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=%2BLHYDJp%2B%2FT1xu%2FesfcvHVUeJweV7MKGjcVE58Twmsk8%3D&amp;reserved=0.
 >>   >>   >>   >>>> 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?.
 >>   >>   >>   >>>>
 >>   >>   >>   >>>> On 11/03/20 8:51 a. m.,
Hamilton Tobon Mosquera via
llvm-dev
 >>   >> wrote:
 >>   >>   >>   >>>>
 >>   >>   >>   >>>> 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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=47sDHCNwfoSt%2B7HDCHg5xe7yTK2IYYPOY6k2NxvGaj8%3D&amp;reserved=0<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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&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.
 >>   >>   >>   >>>>
 >>   >>   >>   >>>>
 >>   >>   >>   >>>> On 10/03/20 7:12 p. m.,
Stefanos Baziotis wrote:
 >>   >>   >>   >>>> 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. :)
 >>   >>   >>   >>>>
 >>   >>   >>   >>>> 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 at lists.llvm.org<mailto:llvm-dev at
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%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=r2Gq2MnAyC8jZzMYPHkEKRnL7p2teev2K3t9%2FdtE9es%3D&amp;reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&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 at
lists.llvm.org<mailto:llvm-dev at lists.llvm.org>
 >>   >>   >>   >>>>
 >>   >>   >>
 >>   >>
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0>
 >>   >>   >>   >>>>
 >>   >>   >>   >>>> 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 at
lists.llvm.org<mailto:llvm-dev at lists.llvm.org>
 >>   >>   >>   >>>>
 >>   >>   >>
 >>   >>
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0
 >>   >>   >>   >>>>
 >>   >>   >>   >>>>
 >>   >>   >>   >>>> 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 at lists.llvm.org
 >>   >>   >>   >>>
 >>   >>   >>
 >>   >>
 >> 
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%3D&amp;reserved=0
 >>   >>   >>   >>>
 >>   >>   >>   >>> 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.

On 28/03/20 11:05 p. m., Doerfert, Johannes wrote:
> On 3/28/20 8:33 PM, Hamilton Tobon Mosquera 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:
>   >> On 3/25/20 5:51 AM, Hamilton Tobon Mosquera 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.
>   >>   >
>   >>   > On 23/03/20 2:21 p. m., Doerfert, Johannes wrote:
>   >>   >>
>   >>   >> On 3/23/20 11:22 AM, Hamilton Tobon Mosquera wrote:
>   >>   >>   > 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
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=yM9OaHxEHZBkU9HNFiMvQFv9zveVQp3kS%2FB4W9qkxKA%3D&reserved=0
>   >> , 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.
>   >>   >>   >
>   >>   >>   > On 23/03/20 11:11 a. m., Doerfert, Johannes
wrote:
>   >>   >>   >> On 3/22/20 12:33 PM, Hamilton Tobon
Mosquera wrote:
>   >>   >>   >>   > 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.
>   >>   >>   >>   >
>   >>   >>   >>   > On 22/03/20 1:26 p. m., Doerfert,
Johannes wrote:
>   >>   >>   >>   >> I did not. The issues sections
on GH was removed.
>   >>   >>   >>   >>
>   >>   >>   >>   >>
________________________________________
>   >>   >>   >>   >> From: Hamilton Tobon Mosquera
<htobonm at eafit.edu.co>
>   >>   >>   >>   >> Sent: Sunday, March 22, 2020
12:21
>   >>   >>   >>   >> To: llvm-dev at 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.
>   >>   >>   >>   >>
>   >>   >>   >>   >> On 21/03/20 10:04 p. m.,
Hamilton Tobon Mosquera via
> llvm-dev
>   >>   >> wrote:
>   >>   >>   >>   >>> Hi there!,
>   >>   >>   >>   >>>
>   >>   >>   >>   >>> Right now I'm starting
with the project proposed by
>   >> Johannes here:
>   >>   >>   >>   >>>
>   >>   >>   >>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=CVDD7HCzFVsPwSpZLOeiO5%2BMNX78sGrDd2eYJ5dWzyE%3D&amp;reserved=0.
>   >>   >>   >>   >>>
>   >>   >>   >>   >>>
>   >>   >>   >>   >>> 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.
>   >>   >>   >>   >>>
>   >>   >>   >>   >>>
>   >>   >>   >>   >>> On 18/03/20 9:21 p. m.,
Doerfert, Johannes wrote:
>   >>   >>   >>   >>>> 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%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&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
>   >>   >>   >>   >>>>
>   >>   >>   >>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=CVDD7HCzFVsPwSpZLOeiO5%2BMNX78sGrDd2eYJ5dWzyE%3D&amp;reserved=0
>   >>   >>   >>   >>>> 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 at anl.gov
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>>
________________________________________
>   >>   >>   >>   >>>> From: Hamilton Tobon
Mosquera <htobonm at eafit.edu.co>
>   >>   >>   >>   >>>> Sent: Wednesday, March
18, 2020 19:14
>   >>   >>   >>   >>>> To: llvm-dev at
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:
>   >>   >>   >>   >>>>
>   >>   >>   >>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD29250&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=lycF%2FQ2WOJtIzRuGEgAu%2BOEw9220%2FnqslSkJMKFENJw%3D&amp;reserved=0.
>   >>   >>   >>   >>>> 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?.
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>> On 11/03/20 8:51 a. m.,
Hamilton Tobon Mosquera via
> llvm-dev
>   >>   >> wrote:
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>> 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%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gAMSEONp%2F8cMDuO0IckxHPXbBnUllJMWVfLHwa6UsBc%3D&amp;reserved=0<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%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&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.
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>> On 10/03/20 7:12 p. m.,
Stefanos Baziotis wrote:
>   >>   >>   >>   >>>> 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. :)
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>> 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 at lists.llvm.org<mailto:llvm-dev at
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%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=Wye5YuFLVUPPESqTRnwsZtimjF0wdUxDetwoCNCib64%3D&amp;reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&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 at
lists.llvm.org<mailto:llvm-dev at lists.llvm.org>
>   >>   >>   >>   >>>>
>   >>   >>   >>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%3D&amp;reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%3D&amp;reserved=0>
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>> 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 at
lists.llvm.org<mailto:llvm-dev at lists.llvm.org>
>   >>   >>   >>   >>>>
>   >>   >>   >>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%3D&amp;reserved=0
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>>
>   >>   >>   >>   >>>> 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 at lists.llvm.org
>   >>   >>   >>   >>>
>   >>   >>   >>
>   >>   >>
>   >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%3D&amp;reserved=0
>   >>   >>   >>   >>>
>   >>   >>   >>   >>> 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.