[llvm-dev] GSoC Interested student

Sat Mar 21 19:04:39 PDT 2020

Hi there!,

Right now I'm starting with the project proposed by Johannes here: 
https://github.com/llvm/llvm-project/issues/180.

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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=jm6nUVSVZtG%2FGG5fIGEPsNIPdrZP1PgFrtGiOyxuWWU%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=YFUCh51fT%2FGq06rf6HxMB6gcpYmVq%2Bds9N%2FeRcB4UkY%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=kWyMuofWqIb0%2BcB67JG05KV3EeDyTDIHueQhSNQt79o%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=KsYBbUlCPFfngpKmHdfgHQkoKV5YOvaCctWCBs87LiE%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=KsYBbUlCPFfngpKmHdfgHQkoKV5YOvaCctWCBs87LiE%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=EmsjHiPFFUrLrc8STVzvUJUWzBJZFQZXkB7lVwOWPpk%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=EmsjHiPFFUrLrc8STVzvUJUWzBJZFQZXkB7lVwOWPpk%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=iyzDSiDzS%2FoOjhkxU09r1XcJ6ghvTZJ0eJSfn3dfyMs%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=iyzDSiDzS%2FoOjhkxU09r1XcJ6ghvTZJ0eJSfn3dfyMs%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%7Cda7fb6508f0347e123f008d7cba3de83%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=iyzDSiDzS%2FoOjhkxU09r1XcJ6ghvTZJ0eJSfn3dfyMs%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.