thr3ads.net - llvm dev - [llvm-dev] GSoC Interested student [Mar 2020]

If this information is useful, please help other people find it:
Share via:

Hamilton Tobon Mosquera via llvm-dev

2020-Mar-25 10:51 UTC

[llvm-dev] GSoC Interested student

Hi Johannes,

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

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

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

Thanks.

Hamilton.

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&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=XIo6OoK0Xnhat9bZbTNz9Rq9hYvBQhknTxULoA8Ma%2Fk%3D&amp;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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=PS4a99t8C7cFxrhlDsCma2nzGHGPdPne8hemx2lr%2Fms%3D&amp;reserved=0
>   >>   >>>> and the OpenMPOpt pass.
>   >>   >>>> There are multiple bigger tasks towards
better offloading, one of
>   >>   >>>> which is described here
>   >>   >>>>
>   >>
>
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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=QNwz3GBwKHxr7bOZurm6rPtD6IwD2XuBGC84dk%2F8CfM%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%3D&amp;reserved=0>
>   >>   >>>>
>   >>   >>>> I have a question. When there is a mismatch
in the types
> (return type
>   >>   >>>> or argument types) between the function
declaration found and the
>   >>   >>>> runtime library function, what should I do?.
Show an error or a
>   >>   >>>> warning?, continuing with the pass or exiting
the program?.
>   >>   >>>>
>   >>   >>>> Thanks.
>   >>   >>>>
>   >>   >>>>
>   >>   >>>> 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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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.

Doerfert, Johannes via llvm-dev

2020-Mar-25 15:41 UTC

head link

[llvm-dev] GSoC Interested student

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&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=XIo6OoK0Xnhat9bZbTNz9Rq9hYvBQhknTxULoA8Ma%2Fk%3D&amp;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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=PS4a99t8C7cFxrhlDsCma2nzGHGPdPne8hemx2lr%2Fms%3D&amp;reserved=0
 >>   >>   >>>> and the OpenMPOpt pass.
 >>   >>   >>>> There are multiple bigger tasks towards
better
offloading, one of
 >>   >>   >>>> which is described here
 >>   >>   >>>>
 >>   >>
 >> 
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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=QNwz3GBwKHxr7bOZurm6rPtD6IwD2XuBGC84dk%2F8CfM%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%3D&amp;reserved=0>
 >>   >>   >>>>
 >>   >>   >>>> I have a question. When there is a
mismatch in the types
 >> (return type
 >>   >>   >>>> or argument types) between the function
declaration
found and the
 >>   >>   >>>> runtime library function, what should I
do?. Show an
error or a
 >>   >>   >>>> warning?, continuing with the pass or
exiting the program?.
 >>   >>   >>>>
 >>   >>   >>>> Thanks.
 >>   >>   >>>>
 >>   >>   >>>>
 >>   >>   >>>> 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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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.
 >

Hamilton Tobon Mosquera via llvm-dev

2020-Mar-29 01:33 UTC

head link

[llvm-dev] 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&amp;data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&amp;sdata=nhtF7ZKL3iPAQEJpIgCFb0cpcPGQYRvOHuPmkxyI6H8%3D&amp;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.

Reasonably Related Threads

Search for more reasonably related threads

llvm dev - Mar 2020 - GSoC Interested student

[llvm-dev] GSoC Interested student

[llvm-dev] GSoC Interested student

[llvm-dev] GSoC Interested student

Reasonably Related Threads