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&data=01%7C01%7Chtobonm%40eafit.edu.co%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=XIo6OoK0Xnhat9bZbTNz9Rq9hYvBQhknTxULoA8Ma%2Fk%3D&reserved=0 , the mapping request, > > which may or may not cause data-transfers, are part of the call to > __tgt_target_teams[_nowait] and __tgt_target_data_update[_nowait]. Start > by reading the impl. of those two functions in the openmp runtime > (llvm-project/openmp/...). > > Let me know if this helps and once you have new questions. > > Cheers, > Johannes > > P.S. I merged your first patch. You didn't run all test before. I did > and noticed that we failed one because of type mismatches. Fixed them > before I merged your patch. > > > > > #include <omp.h> > > #include <stdio.h> > > > > #define size 1000000 > > double a[size]; > > > > int main() { > > > > #pragma omp target map(to: a[0:size]) > > #pragma omp teams distribute > > for (int i = 0; i < size; ++i) { > > a[i] = a[i]*(i*i)/2; > > } > > > > return 0; > > } > > > > > > PD: The command to compile is this: > > clang -v -save-temps \ > > -I/path/to/llvm/release/9.x/include -fopenmp \ > > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \ > > -march=sm_61 memory_transfer.c -o memory_transfer > > > > Thanks. > > Hamilton. > > > > On 23/03/20 11:11 a. m., Doerfert, Johannes wrote: > >> On 3/22/20 12:33 PM, Hamilton Tobon Mosquera wrote: > >> > Ahhh I understand. So I believe the project you posted is still > >> opened. Do you any advice on the question I posted before?. > >> > >> Yes. Still open. You should use this as the main part of the GSoC > >> proposal (if you are interested). Can you repeat the question? > >> > >> > > >> > Thank you. > >> > > >> > On 22/03/20 1:26 p. m., Doerfert, Johannes wrote: > >> >> I did not. The issues sections on GH was removed. > >> >> > >> >> ________________________________________ > >> >> From: Hamilton Tobon Mosquera <htobonm at eafit.edu.co> > >> >> Sent: Sunday, March 22, 2020 12:21 > >> >> To: llvm-dev at lists.llvm.org > >> >> Cc: Doerfert, Johannes > >> >> Subject: Re: [llvm-dev] GSoC Interested student > >> >> > >> >> Hi there Johannes!, > >> >> > >> >> I was checking out the issue (about latency hiding when > transferring > >> >> memory from host to device) you posted about a week ago, but > it's not > >> >> there anymore. Did you remove it?. > >> >> > >> >> Thanks. > >> >> > >> >> Hamilton. > >> >> > >> >> On 21/03/20 10:04 p. m., Hamilton Tobon Mosquera via llvm-dev > wrote: > >> >>> Hi there!, > >> >>> > >> >>> Right now I'm starting with the project proposed by Johannes here: > >> >>> > >> > https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&data=01%7C01%7Chtobonm%40eafit.edu.co%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=PS4a99t8C7cFxrhlDsCma2nzGHGPdPne8hemx2lr%2Fms%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=QNwz3GBwKHxr7bOZurm6rPtD6IwD2XuBGC84dk%2F8CfM%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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. > > > > > La información contenida en este correo electrónico está dirigida únicamente a su destinatario y puede contener información confidencial, material privilegiado o información protegida por derecho de autor. Está prohibida cualquier copia, utilización, indebida retención, modificación, difusión, distribución o reproducción total o parcial. Si usted recibe este mensaje por error, por favor contacte al remitente y elimínelo. La información aquí contenida es responsabilidad exclusiva de su remitente por lo tanto la Universidad EAFIT no se hace responsable de lo que el mensaje contenga. The information contained in this email is addressed to its recipient only and may contain confidential information, privileged material or information protected by copyright. Its prohibited any copy, use, improper retention, modification, dissemination, distribution or total or partial reproduction. If you receive this message by error, please contact the sender and delete it. The information contained herein is the sole responsibility of the sender therefore Universidad EAFIT is not responsible for what the message contains.
On 3/25/20 5:51 AM, Hamilton Tobon Mosquera wrote:
> Hi Johannes,
>
> Sorry about the failing tests. I completely forgot about running all
> of them, I only ran the tests for the openmp ipo transforms.
No worries, it happens.
> I found the GH issue in cache. I can post it here if you want.
Yes, please.
> I analyzed these runtime calls, __tgt_target_teams_[nowait]. What I
> don't understand about the problem we want to solve (splitting the
> memory transfer into issue and wait and legally moving code inside) is
> that the memory transfer is made inside those runtime calls.
> Therefore, the runtime calls to the memory transfer cannot be analyzed
> in the IR. Could you please explain a bit further if the memory
> transfer should split in the IR? or split __tgt_target_teams_[nowait]
> in such a way that the memory transfer is made in the first part, and
> the wait in the second part.
Exactly what you said in the last sentence. Split the runtime call into
two, issue and wait. Then we move issue "up" the execution path and
wait
"down".
Cheers,
Johannes
> Thanks.
>
> Hamilton.
>
> On 23/03/20 2:21 p. m., Doerfert, Johannes wrote:
>>
>> On 3/23/20 11:22 AM, Hamilton Tobon Mosquera wrote:
>> > Hi Johannes,
>> >
>> > Yes, I'm highly interested, I already started. It would be
nice
if you
>> > post the details again. I didn't take note of them, I was
always
>> > looking at the issue in GH.
>>
>> I asked if we can make the readonly but unsure yet. I can write it
again
>> but it'll take ma some time.
>>
>>
>> > The problem I'm facing right now is this:
>> > The first task is identifying OpenMP API calls referred to
device memory
>> > allocation. I first understood how the compiler driver
orchestrates the
>> > process of offloading OpenMP regions to devices, and now
I'm
looking at
>> > the step of the pipeline that generates LLVM IR for both host
and
>> > device. I wrote a simple C program that uses OpenMP to map a
local array
>> > to the device main memory. Then I'm looking at the device
and
host LLVM
>> > IR representation of this program, but I can't find any
runtime call
>> > referred to device memory allocation. In the host IR I see that
the
>> > array is passed to the outlined function and stored in a
register inside
>> > the outlined function, but there's no memory allocation
runtime
call.
>> > And in the device IR the array is passed to the main offloading
>> > function.
>> > Any suggestion or documentation on how that memory transfer is
>> > achieved?, and how to track its API calls. It might be that
I'm not
>> > using the proper omp directive, so here is C code:
>>
>> Take a look at
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=XIo6OoK0Xnhat9bZbTNz9Rq9hYvBQhknTxULoA8Ma%2Fk%3D&reserved=0
, the mapping request,
>>
>> which may or may not cause data-transfers, are part of the call to
>> __tgt_target_teams[_nowait] and __tgt_target_data_update[_nowait].
Start
>> by reading the impl. of those two functions in the openmp runtime
>> (llvm-project/openmp/...).
>>
>> Let me know if this helps and once you have new questions.
>>
>> Cheers,
>> Johannes
>>
>> P.S. I merged your first patch. You didn't run all test before. I
did
>> and noticed that we failed one because of type mismatches. Fixed them
>> before I merged your patch.
>>
>> >
>> > #include <omp.h>
>> > #include <stdio.h>
>> >
>> > #define size 1000000
>> > double a[size];
>> >
>> > int main() {
>> >
>> > #pragma omp target map(to: a[0:size])
>> > #pragma omp teams distribute
>> > for (int i = 0; i < size; ++i) {
>> > a[i] = a[i]*(i*i)/2;
>> > }
>> >
>> > return 0;
>> > }
>> >
>> >
>> > PD: The command to compile is this:
>> > clang -v -save-temps \
>> > -I/path/to/llvm/release/9.x/include -fopenmp \
>> > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \
>> > -march=sm_61 memory_transfer.c -o memory_transfer
>> >
>> > Thanks.
>> > Hamilton.
>> >
>> > On 23/03/20 11:11 a. m., Doerfert, Johannes wrote:
>> >> On 3/22/20 12:33 PM, Hamilton Tobon Mosquera wrote:
>> >> > Ahhh I understand. So I believe the project you
posted is still
>> >> opened. Do you any advice on the question I posted before?.
>> >>
>> >> Yes. Still open. You should use this as the main part of
the GSoC
>> >> proposal (if you are interested). Can you repeat the
question?
>> >>
>> >> >
>> >> > Thank you.
>> >> >
>> >> > On 22/03/20 1:26 p. m., Doerfert, Johannes wrote:
>> >> >> I did not. The issues sections on GH was
removed.
>> >> >>
>> >> >> ________________________________________
>> >> >> From: Hamilton Tobon Mosquera <htobonm at
eafit.edu.co>
>> >> >> Sent: Sunday, March 22, 2020 12:21
>> >> >> To: llvm-dev at lists.llvm.org
>> >> >> Cc: Doerfert, Johannes
>> >> >> Subject: Re: [llvm-dev] GSoC Interested student
>> >> >>
>> >> >> Hi there Johannes!,
>> >> >>
>> >> >> I was checking out the issue (about latency
hiding when
>> transferring
>> >> >> memory from host to device) you posted about a
week ago, but
>> it's not
>> >> >> there anymore. Did you remove it?.
>> >> >>
>> >> >> Thanks.
>> >> >>
>> >> >> Hamilton.
>> >> >>
>> >> >> On 21/03/20 10:04 p. m., Hamilton Tobon Mosquera
via llvm-dev
>> wrote:
>> >> >>> Hi there!,
>> >> >>>
>> >> >>> Right now I'm starting with the project
proposed by
Johannes here:
>> >> >>>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&data=01%7C01%7Chtobonm%40eafit.edu.co%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=PS4a99t8C7cFxrhlDsCma2nzGHGPdPne8hemx2lr%2Fms%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=cGz5OkdMjvpWRtIkIjkKLXHj57uzPlu7BRb62eHGiJI%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=QNwz3GBwKHxr7bOZurm6rPtD6IwD2XuBGC84dk%2F8CfM%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pVa4w%2FncMMrFWSf0oKf0aoRdBXoVkjm42qOucKqLfP0%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=BDpqnRQqD4sYU5rFbka4n8rFQcF0QkH2FbZU1ftB6bQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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%7C1c12e490f3804f6a772d08d7cf570a3a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gRHzlUxQPBoKCqNLJpCNLTxxGKmW1blBZjqanouEolQ%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.
>> >
>>
>>
>> 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
[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&data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=nhtF7ZKL3iPAQEJpIgCFb0cpcPGQYRvOHuPmkxyI6H8%3D&reserved=0 > , the mapping request, > >> > >> which may or may not cause data-transfers, are part of the call to > >> __tgt_target_teams[_nowait] and __tgt_target_data_update[_nowait]. Start > >> by reading the impl. of those two functions in the openmp runtime > >> (llvm-project/openmp/...). > >> > >> Let me know if this helps and once you have new questions. > >> > >> Cheers, > >> Johannes > >> > >> P.S. I merged your first patch. You didn't run all test before. I did > >> and noticed that we failed one because of type mismatches. Fixed them > >> before I merged your patch. > >> > >> > > >> > #include <omp.h> > >> > #include <stdio.h> > >> > > >> > #define size 1000000 > >> > double a[size]; > >> > > >> > int main() { > >> > > >> > #pragma omp target map(to: a[0:size]) > >> > #pragma omp teams distribute > >> > for (int i = 0; i < size; ++i) { > >> > a[i] = a[i]*(i*i)/2; > >> > } > >> > > >> > return 0; > >> > } > >> > > >> > > >> > PD: The command to compile is this: > >> > clang -v -save-temps \ > >> > -I/path/to/llvm/release/9.x/include -fopenmp \ > >> > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \ > >> > -march=sm_61 memory_transfer.c -o memory_transfer > >> > > >> > Thanks. > >> > Hamilton. > >> > > >> > On 23/03/20 11:11 a. m., Doerfert, Johannes wrote: > >> >> On 3/22/20 12:33 PM, Hamilton Tobon Mosquera wrote: > >> >> > Ahhh I understand. So I believe the project you posted is still > >> >> opened. Do you any advice on the question I posted before?. > >> >> > >> >> Yes. Still open. You should use this as the main part of the GSoC > >> >> proposal (if you are interested). Can you repeat the question? > >> >> > >> >> > > >> >> > Thank you. > >> >> > > >> >> > On 22/03/20 1:26 p. m., Doerfert, Johannes wrote: > >> >> >> I did not. The issues sections on GH was removed. > >> >> >> > >> >> >> ________________________________________ > >> >> >> From: Hamilton Tobon Mosquera <htobonm at eafit.edu.co> > >> >> >> Sent: Sunday, March 22, 2020 12:21 > >> >> >> To: llvm-dev at lists.llvm.org > >> >> >> Cc: Doerfert, Johannes > >> >> >> Subject: Re: [llvm-dev] GSoC Interested student > >> >> >> > >> >> >> Hi there Johannes!, > >> >> >> > >> >> >> I was checking out the issue (about latency hiding when > >> transferring > >> >> >> memory from host to device) you posted about a week ago, but > >> it's not > >> >> >> there anymore. Did you remove it?. > >> >> >> > >> >> >> Thanks. > >> >> >> > >> >> >> Hamilton. > >> >> >> > >> >> >> On 21/03/20 10:04 p. m., Hamilton Tobon Mosquera via llvm-dev > >> wrote: > >> >> >>> Hi there!, > >> >> >>> > >> >> >>> Right now I'm starting with the project proposed by > Johannes here: > >> >> >>> > >> >> > >> > https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&data=01%7C01%7Chtobonm%40eafit.edu.co%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=wyVbkLe8jhPzEIH6zHRwZZLdPHCKXPuXQNW%2BOhMHxv4%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=zQQajI2aySpFNO4EieIW8u5I2ei5i6myfdescuN0%2FtY%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=wyVbkLe8jhPzEIH6zHRwZZLdPHCKXPuXQNW%2BOhMHxv4%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=%2BLHYDJp%2B%2FT1xu%2FesfcvHVUeJweV7MKGjcVE58Twmsk8%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=47sDHCNwfoSt%2B7HDCHg5xe7yTK2IYYPOY6k2NxvGaj8%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=47sDHCNwfoSt%2B7HDCHg5xe7yTK2IYYPOY6k2NxvGaj8%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=r2Gq2MnAyC8jZzMYPHkEKRnL7p2teev2K3t9%2FdtE9es%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=r2Gq2MnAyC8jZzMYPHkEKRnL7p2teev2K3t9%2FdtE9es%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%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%7Cd2a1829610714bf161b508d7d0d2f5ae%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=3gX0CioY5Qx7XMxrCbXLLgwBURQhx125bNteDljB0lY%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. > >> > > >> > >> > >> 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.