[llvm-dev] GSoC Interested student
Hamilton Tobon Mosquera via llvm-dev
llvm-dev at lists.llvm.org
Mon Mar 23 09:22:55 PDT 2020
Hi Johannes,
Yes, I'm highly interested, I already started. It would be nice if you post the details again. I didn't take note of them, I was always looking at the issue in GH.
The problem I'm facing right now is this:
The first task is identifying OpenMP API calls referred to device memory
allocation. I first understood how the compiler driver orchestrates the
process of offloading OpenMP regions to devices, and now I'm looking at
the step of the pipeline that generates LLVM IR for both host and
device. I wrote a simple C program that uses OpenMP to map a local array
to the device main memory. Then I'm looking at the device and host LLVM
IR representation of this program, but I can't find any runtime call
referred to device memory allocation. In the host IR I see that the
array is passed to the outlined function and stored in a register inside
the outlined function, but there's no memory allocation runtime call.
And in the device IR the array is passed to the main offloading
function.
Any suggestion or documentation on how that memory transfer is
achieved?, and how to track its API calls. It might be that I'm not
using the proper omp directive, so here is C code:
#include <omp.h>
#include <stdio.h>
#define size 1000000
double a[size];
int main() {
#pragma omp target map(to: a[0:size])
#pragma omp teams distribute
for (int i = 0; i < size; ++i) {
a[i] = a[i]*(i*i)/2;
}
return 0;
}
PD: The command to compile is this:
clang -v -save-temps \
-I/path/to/llvm/release/9.x/include -fopenmp \
-fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
-march=sm_61 memory_transfer.c -o memory_transfer
Thanks.
Hamilton.
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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=QU6c8JHe33fkyAhEKhj3%2BBQk3AiAf7BhbXhB7ny2%2Fxk%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=6nejCYt05TQdtstWkksVMf2%2FZyarC3GLEifC0Iyw%2Bic%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=QU6c8JHe33fkyAhEKhj3%2BBQk3AiAf7BhbXhB7ny2%2Fxk%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=vPF1RTXqSVctiDbLZXXqKL3HPCWWjrbMPEkAAcgq1QQ%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=SbO3ePJdrdDaA1RjUPgOzS7vplRKOe8ZjY68DeM66NQ%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=SbO3ePJdrdDaA1RjUPgOzS7vplRKOe8ZjY68DeM66NQ%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=UbADZ7YRMOi3%2FCVt%2FGvUMYCTw%2Bgmr8psk2QPVEOC8CE%3D&reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fbugs.llvm.org%2Fshow_bug.cgi%3Fid%3D40253&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=UbADZ7YRMOi3%2FCVt%2FGvUMYCTw%2Bgmr8psk2QPVEOC8CE%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=auXp2U8Qmo8DPGkDEn8k5NwVGfAHZUTKTlhxPcBP%2FrY%3D&reserved=0<https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-dev&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=auXp2U8Qmo8DPGkDEn8k5NwVGfAHZUTKTlhxPcBP%2FrY%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=auXp2U8Qmo8DPGkDEn8k5NwVGfAHZUTKTlhxPcBP%2FrY%3D&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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C27650f5cf4e24c5920be08d7cf3c8209%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=auXp2U8Qmo8DPGkDEn8k5NwVGfAHZUTKTlhxPcBP%2FrY%3D&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.
More information about the llvm-dev
mailing list