Hamilton Tobon Mosquera via llvm-dev
2020-Apr-07 12:32 UTC
[llvm-dev] GSoC Interested student
Hi Johannes, In order to create the testcases with the FIXMEs showing how it should look after the transformation, I've been trying to come up with a way of splitting the runtime call __tgt_target_teams based on Silei's recent patch D77005. I like the idea of using its asynchronous version __tgt_target_teams_nowait, but I think we would have to modify it, either to return the async_info struct: __tgt_async_info __tgt_target_teams_nowait(...) // This would break the current canonical way of returning a status code. or to receive it as a pointer and modify its content, then waiting on that async_info object: int __tgt_target_teams_nowait(..., __tgt_async_info *async_info = nullptr) The first way leads to modify the Clang's code generation, and I'm not sure if that is a good idea. Concerning the second way, I'm not sure if having a default argument would lead also to modify the Clang's code generation. Another possible solution would be to implement another runtime function, avoiding the need to modify the _nowait version. This would lead to implement new runtime functions for each runtime call that implies data transfers, which is perhaps worse. I don't know which approach is better, should I look for a way of modifying Clang's code generation in case of needed?. I'm trying to be as less disruptive as possible. Can you give me some advice?. Thank you. Hamilton. On 28/03/20 11:05 p. m., Doerfert, Johannes wrote:> On 3/28/20 8:33 PM, Hamilton Tobon Mosquera wrote: > > Hi Johannes, > > > > I've been thinking on how to implement a solution for the problem and > > I have some questions: > > > > 1. You mention in the problem description that the current > > asynchronous functions can be used, and probably modified. The problem > > is that those runtime functions with _nowait at the end are not > > asynchronous. Looking at their implementations, they just wait for > > other tasks is there are pending tasks, and then call its synchronous > > versions. So, is it a good idea to wrap them in functions that call > > them as asynchronous (using std::async), then return an object to wait > > on?. Then, waiting on the object returned. > > Take a look at D77005. Shilei introduced some real askync versions for a > different reason already. You can build on top of that patch and extend > it if needed further. > > > > 2. I'm thinking on replacing the data transfer inside > > __tgt_target_teams with __tgt_target_data_begin. The problem with this > > is that __tgt_target_data_begin does not initialize [first]private > > variables, and doesn't check for lambda mappings. Therefore, my idea > > to go around this problem is to create another wrapper function F, > > that calls __tgt_target_data_begin, and executes the necessary code > > for lambda mappings check and [first]private variables, and finally, > > calling F as the asynchronous function. What is good about this, is > > that the code concerning data transfer inside __tgt_target_teams can > > be replaced with F. > > Interesting. I'll need to see small examples. For now, focus on the > simple case, we'll get to the rest later. However, we should really > start by creating test cases that show the initial code and FIXMEs that > describe how we want it to look like. > > > > 3. Also, I was wondering if I must to be aware of (split them into > > issue and wait) all the other constructs that imply data transfers > > (update, data enter/exit), or just focus on __tgt_target[_nowait] and > > __tgt_target_teams[_nowait]. > > We want to eventually split all transfers into two but let's start with > a single one. > > > > 4. This might be a dumb question, but just to make sure. Do I have to > > move the wait call into the offloaded region just before the data is > > used (because I think this would imply analyzing the offloaded > > bitcode, which is in another module)?, or move it just before the > > kernel is executed?. > > We split, then we move them apart. The cannot be moved over instructions > that might change behavior if we move them, e.g. an aliasing memory > operation or a kernel call that uses the data. > > > > 5. Do you have a template for the GSoC proposal. That is, is it plain > > text or pdf or ...?. Also, do you have sort of a structure that we > > have to follow?, for example, Abstract -> Benefits -> Deliverables > > ..., or do we have freedom on that?. Your post in hte GSoC page only > > mentions that we must include our background, but nothing about the > > format of the proposal. > > I do not have one. I thought they might or they have an FAQ that > describes what to do. I know they have some required things, so check > their webpage. I think you have to use google docs, I might be wrong. > Your categories look good, make the deliverables not too specific but > describe more ideas to work on. We'll make tasks as we go as you are not > the only one working on this general area. > > Feel free to send me a draft of the proposal for feedback. > > Cheers, > Johannes > > > Thank you. > > > > Hamilton. > > > > > > On``` 25/03/20 11:41 a. m., Doerfert, Johannes wrote: > >> On 3/25/20 5:51 AM, Hamilton Tobon Mosquera wrote: > >> > Hi Johannes, > >> > > >> > Sorry about the failing tests. I completely forgot about running all > >> > of them, I only ran the tests for the openmp ipo transforms. > >> > >> No worries, it happens. > >> > >> > >> > I found the GH issue in cache. I can post it here if you want. > >> > >> Yes, please. > >> > >> > >> > I analyzed these runtime calls, __tgt_target_teams_[nowait]. What I > >> > don't understand about the problem we want to solve (splitting the > >> > memory transfer into issue and wait and legally moving code > inside) is > >> > that the memory transfer is made inside those runtime calls. > >> > Therefore, the runtime calls to the memory transfer cannot be > analyzed > >> > in the IR. Could you please explain a bit further if the memory > >> > transfer should split in the IR? or split > __tgt_target_teams_[nowait] > >> > in such a way that the memory transfer is made in the first > part, and > >> > the wait in the second part. > >> > >> Exactly what you said in the last sentence. Split the runtime call into > >> two, issue and wait. Then we move issue "up" the execution path and wait > >> "down". > >> > >> Cheers, > >> Johannes > >> > >> > >> > Thanks. > >> > > >> > Hamilton. > >> > > >> > On 23/03/20 2:21 p. m., Doerfert, Johannes wrote: > >> >> > >> >> On 3/23/20 11:22 AM, Hamilton Tobon Mosquera wrote: > >> >> > Hi Johannes, > >> >> > > >> >> > Yes, I'm highly interested, I already started. It would be nice > >> if you > >> >> > post the details again. I didn't take note of them, I was > always > >> >> > looking at the issue in GH. > >> >> > >> >> I asked if we can make the readonly but unsure yet. I can write > it again > >> >> but it'll take ma some time. > >> >> > >> >> > >> >> > The problem I'm facing right now is this: > >> >> > The first task is identifying OpenMP API calls referred to > >> device memory > >> >> > allocation. I first understood how the compiler driver > >> orchestrates the > >> >> > process of offloading OpenMP regions to devices, and now I'm > >> looking at > >> >> > the step of the pipeline that generates LLVM IR for both > host and > >> >> > device. I wrote a simple C program that uses OpenMP to map a > >> local array > >> >> > to the device main memory. Then I'm looking at the device and > >> host LLVM > >> >> > IR representation of this program, but I can't find any > runtime call > >> >> > referred to device memory allocation. In the host IR I see > that the > >> >> > array is passed to the outlined function and stored in a > >> register inside > >> >> > the outlined function, but there's no memory allocation runtime > >> call. > >> >> > And in the device IR the array is passed to the main offloading > >> >> > function. > >> >> > Any suggestion or documentation on how that memory transfer is > >> >> > achieved?, and how to track its API calls. It might be that > I'm not > >> >> > using the proper omp directive, so here is C code: > >> >> > >> >> Take a look at > >> > https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=yM9OaHxEHZBkU9HNFiMvQFv9zveVQp3kS%2FB4W9qkxKA%3D&reserved=0 > >> , the mapping request, > >> >> > >> >> which may or may not cause data-transfers, are part of the call to > >> >> __tgt_target_teams[_nowait] and > __tgt_target_data_update[_nowait]. Start > >> >> by reading the impl. of those two functions in the openmp runtime > >> >> (llvm-project/openmp/...). > >> >> > >> >> Let me know if this helps and once you have new questions. > >> >> > >> >> Cheers, > >> >> Johannes > >> >> > >> >> P.S. I merged your first patch. You didn't run all test before. > I did > >> >> and noticed that we failed one because of type mismatches. > Fixed them > >> >> before I merged your patch. > >> >> > >> >> > > >> >> > #include <omp.h> > >> >> > #include <stdio.h> > >> >> > > >> >> > #define size 1000000 > >> >> > double a[size]; > >> >> > > >> >> > int main() { > >> >> > > >> >> > #pragma omp target map(to: a[0:size]) > >> >> > #pragma omp teams distribute > >> >> > for (int i = 0; i < size; ++i) { > >> >> > a[i] = a[i]*(i*i)/2; > >> >> > } > >> >> > > >> >> > return 0; > >> >> > } > >> >> > > >> >> > > >> >> > PD: The command to compile is this: > >> >> > clang -v -save-temps \ > >> >> > -I/path/to/llvm/release/9.x/include -fopenmp \ > >> >> > -fopenmp-targets=nvptx64-nvidia-cuda -Xopenmp-target \ > >> >> > -march=sm_61 memory_transfer.c -o memory_transfer > >> >> > > >> >> > Thanks. > >> >> > Hamilton. > >> >> > > >> >> > On 23/03/20 11:11 a. m., Doerfert, Johannes wrote: > >> >> >> On 3/22/20 12:33 PM, Hamilton Tobon Mosquera wrote: > >> >> >> > Ahhh I understand. So I believe the project you posted > is still > >> >> >> opened. Do you any advice on the question I posted before?. > >> >> >> > >> >> >> Yes. Still open. You should use this as the main part of > the GSoC > >> >> >> proposal (if you are interested). Can you repeat the question? > >> >> >> > >> >> >> > > >> >> >> > Thank you. > >> >> >> > > >> >> >> > On 22/03/20 1:26 p. m., Doerfert, Johannes wrote: > >> >> >> >> I did not. The issues sections on GH was removed. > >> >> >> >> > >> >> >> >> ________________________________________ > >> >> >> >> From: Hamilton Tobon Mosquera <htobonm at eafit.edu.co> > >> >> >> >> Sent: Sunday, March 22, 2020 12:21 > >> >> >> >> To: llvm-dev at lists.llvm.org > >> >> >> >> Cc: Doerfert, Johannes > >> >> >> >> Subject: Re: [llvm-dev] GSoC Interested student > >> >> >> >> > >> >> >> >> Hi there Johannes!, > >> >> >> >> > >> >> >> >> I was checking out the issue (about latency hiding when > >> >> transferring > >> >> >> >> memory from host to device) you posted about a week > ago, but > >> >> it's not > >> >> >> >> there anymore. Did you remove it?. > >> >> >> >> > >> >> >> >> Thanks. > >> >> >> >> > >> >> >> >> Hamilton. > >> >> >> >> > >> >> >> >> On 21/03/20 10:04 p. m., Hamilton Tobon Mosquera via > llvm-dev > >> >> wrote: > >> >> >> >>> Hi there!, > >> >> >> >>> > >> >> >> >>> Right now I'm starting with the project proposed by > >> Johannes here: > >> >> >> >>> > >> >> >> > >> >> > >> > https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=CVDD7HCzFVsPwSpZLOeiO5%2BMNX78sGrDd2eYJ5dWzyE%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=hN9WGEO%2F3quio7A9tvtIDHxznURuITYVBszFSmqM5M8%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=CVDD7HCzFVsPwSpZLOeiO5%2BMNX78sGrDd2eYJ5dWzyE%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=lycF%2FQ2WOJtIzRuGEgAu%2BOEw9220%2FnqslSkJMKFENJw%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gAMSEONp%2F8cMDuO0IckxHPXbBnUllJMWVfLHwa6UsBc%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gAMSEONp%2F8cMDuO0IckxHPXbBnUllJMWVfLHwa6UsBc%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=Wye5YuFLVUPPESqTRnwsZtimjF0wdUxDetwoCNCib64%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=Wye5YuFLVUPPESqTRnwsZtimjF0wdUxDetwoCNCib64%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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. > > > > > 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 4/7/20 7:32 AM, Hamilton Tobon Mosquera wrote:
> Hi Johannes,
>
> In order to create the testcases with the FIXMEs showing how it
should look after the transformation, I've been trying to come up with a
way of splitting the runtime call __tgt_target_teams based on Silei's
recent patch D77005. I like the idea of using its asynchronous version
__tgt_target_teams_nowait, but I think we would have to modify it,
either to return the async_info struct:
>
> __tgt_async_info __tgt_target_teams_nowait(...) // This would break
the current canonical way of returning a status code.
>
> or to receive it as a pointer and modify its content, then waiting on
that async_info object:
>
> int __tgt_target_teams_nowait(..., __tgt_async_info *async_info =
nullptr)
>
> The first way leads to modify the Clang's code generation, and I'm
not sure if that is a good idea. Concerning the second way, I'm not sure
if having a default argument would lead also to modify the Clang's code
generation.
>
> Another possible solution would be to implement another runtime
function, avoiding the need to modify the _nowait version. This would
lead to implement new runtime functions for each runtime call that
implies data transfers, which is perhaps worse. I don't know which
approach is better, should I look for a way of modifying Clang's code
generation in case of needed?. I'm trying to be as less disruptive as
possible. Can you give me some advice?.
We are in the process of moving clangs OpenMP code generation into the
OpenMPIRBuilder. While doing so we refactor/replace/rewrite stuff as we
see fit. Thus, changing Clang, or better writing a different code
generation in the OpenMPIRBuilder, is always an option.
That said, I was expecting to create two new runtime calls per existing
one but I have not really thought this through. What I imagined is
something vaguely looking like this:
Current:
int omp_transfer_data(device, src, dst, size);
New:
```
// In the runtime {
handle omp_transfer_data_issue(device, src, dst, size) {
// some logic that was in omp_transfer_data before creating and using
// handle
return some_handle; // probably extended __tgt_async_info, can be an
// argument as well.
}
int omp_transfer_data_wait(device, src, dst, size, handle) {
// some logic that was in omp_transfer_data before using handle
return error_code;
}
// }
// In the runtime or in the header
int omp_transfer_data(device, src, dst, size) {
handle = omp_transfer_data_issue(...);
return omp_transfer_data_wait(..., handle);
}
```
Now we can inline omp_transfer_data or replace it by the two calls to
make the issue and wait explicit.
Long story short, we will most certainly need new functions, how they
look is up to you initially. Look at the other functions for inspiration
and if something is off we'll let you know during code review ;)
Does that help? Feel free to ask more questions!
Cheers,
Johannes
P.S. Tomorrow there is an OpenMP in LLVM bi-weekly call which might be
interesting (though it might be too early):
http://lists.llvm.org/pipermail/openmp-dev/2020-April/003292.html
> Thank you.
>
> Hamilton.
>
> On 28/03/20 11:05 p. m., Doerfert, Johannes wrote:
>> On 3/28/20 8:33 PM, Hamilton Tobon Mosquera wrote:
>> > Hi Johannes,
>> >
>> > I've been thinking on how to implement a solution for the
problem and
>> > I have some questions:
>> >
>> > 1. You mention in the problem description that the current
>> > asynchronous functions can be used, and probably modified. The
problem
>> > is that those runtime functions with _nowait at the end are not
>> > asynchronous. Looking at their implementations, they just wait
for
>> > other tasks is there are pending tasks, and then call its
synchronous
>> > versions. So, is it a good idea to wrap them in functions that
call
>> > them as asynchronous (using std::async), then return an object
to wait
>> > on?. Then, waiting on the object returned.
>>
>> Take a look at D77005. Shilei introduced some real askync versions for
a
>> different reason already. You can build on top of that patch and
extend
>> it if needed further.
>>
>>
>> > 2. I'm thinking on replacing the data transfer inside
>> > __tgt_target_teams with __tgt_target_data_begin. The problem
with this
>> > is that __tgt_target_data_begin does not initialize
[first]private
>> > variables, and doesn't check for lambda mappings.
Therefore, my idea
>> > to go around this problem is to create another wrapper function
F,
>> > that calls __tgt_target_data_begin, and executes the necessary
code
>> > for lambda mappings check and [first]private variables, and
finally,
>> > calling F as the asynchronous function. What is good about
this, is
>> > that the code concerning data transfer inside
__tgt_target_teams can
>> > be replaced with F.
>>
>> Interesting. I'll need to see small examples. For now, focus on
the
>> simple case, we'll get to the rest later. However, we should
really
>> start by creating test cases that show the initial code and FIXMEs
that
>> describe how we want it to look like.
>>
>>
>> > 3. Also, I was wondering if I must to be aware of (split them
into
>> > issue and wait) all the other constructs that imply data
transfers
>> > (update, data enter/exit), or just focus on
__tgt_target[_nowait] and
>> > __tgt_target_teams[_nowait].
>>
>> We want to eventually split all transfers into two but let's start
with
>> a single one.
>>
>>
>> > 4. This might be a dumb question, but just to make sure. Do I
have to
>> > move the wait call into the offloaded region just before the
data is
>> > used (because I think this would imply analyzing the offloaded
>> > bitcode, which is in another module)?, or move it just before
the
>> > kernel is executed?.
>>
>> We split, then we move them apart. The cannot be moved over
instructions
>> that might change behavior if we move them, e.g. an aliasing memory
>> operation or a kernel call that uses the data.
>>
>>
>> > 5. Do you have a template for the GSoC proposal. That is, is it
plain
>> > text or pdf or ...?. Also, do you have sort of a structure that
we
>> > have to follow?, for example, Abstract -> Benefits ->
Deliverables
>> > ..., or do we have freedom on that?. Your post in hte GSoC page
only
>> > mentions that we must include our background, but nothing about
the
>> > format of the proposal.
>>
>> I do not have one. I thought they might or they have an FAQ that
>> describes what to do. I know they have some required things, so check
>> their webpage. I think you have to use google docs, I might be wrong.
>> Your categories look good, make the deliverables not too specific but
>> describe more ideas to work on. We'll make tasks as we go as you
are not
>> the only one working on this general area.
>>
>> Feel free to send me a draft of the proposal for feedback.
>>
>> Cheers,
>> Johannes
>>
>> > Thank you.
>> >
>> > Hamilton.
>> >
>> >
>> > On``` 25/03/20 11:41 a. m., Doerfert, Johannes wrote:
>> >> On 3/25/20 5:51 AM, Hamilton Tobon Mosquera wrote:
>> >> > Hi Johannes,
>> >> >
>> >> > Sorry about the failing tests. I completely forgot
about
running all
>> >> > of them, I only ran the tests for the openmp ipo
transforms.
>> >>
>> >> No worries, it happens.
>> >>
>> >>
>> >> > I found the GH issue in cache. I can post it here if
you want.
>> >>
>> >> Yes, please.
>> >>
>> >>
>> >> > I analyzed these runtime calls,
__tgt_target_teams_[nowait]. What I
>> >> > don't understand about the problem we want to
solve
(splitting the
>> >> > memory transfer into issue and wait and legally
moving code
>> inside) is
>> >> > that the memory transfer is made inside those
runtime calls.
>> >> > Therefore, the runtime calls to the memory transfer
cannot be
>> analyzed
>> >> > in the IR. Could you please explain a bit further if
the memory
>> >> > transfer should split in the IR? or split
>> __tgt_target_teams_[nowait]
>> >> > in such a way that the memory transfer is made in
the first
>> part, and
>> >> > the wait in the second part.
>> >>
>> >> Exactly what you said in the last sentence. Split the
runtime
call into
>> >> two, issue and wait. Then we move issue "up" the
execution path
and wait
>> >> "down".
>> >>
>> >> Cheers,
>> >> Johannes
>> >>
>> >>
>> >> > Thanks.
>> >> >
>> >> > Hamilton.
>> >> >
>> >> > On 23/03/20 2:21 p. m., Doerfert, Johannes wrote:
>> >> >>
>> >> >> On 3/23/20 11:22 AM, Hamilton Tobon Mosquera
wrote:
>> >> >> > Hi Johannes,
>> >> >> >
>> >> >> > Yes, I'm highly interested, I already
started. It
would be nice
>> >> if you
>> >> >> > post the details again. I didn't take
note of them, I was
>> always
>> >> >> > looking at the issue in GH.
>> >> >>
>> >> >> I asked if we can make the readonly but unsure
yet. I can
write
>> it again
>> >> >> but it'll take ma some time.
>> >> >>
>> >> >>
>> >> >> > The problem I'm facing right now is
this:
>> >> >> > The first task is identifying OpenMP API
calls referred to
>> >> device memory
>> >> >> > allocation. I first understood how the
compiler driver
>> >> orchestrates the
>> >> >> > process of offloading OpenMP regions to
devices, and
now I'm
>> >> looking at
>> >> >> > the step of the pipeline that generates
LLVM IR for both
>> host and
>> >> >> > device. I wrote a simple C program that
uses OpenMP to
map a
>> >> local array
>> >> >> > to the device main memory. Then I'm
looking at the
device and
>> >> host LLVM
>> >> >> > IR representation of this program, but I
can't find any
>> runtime call
>> >> >> > referred to device memory allocation. In
the host IR I see
>> that the
>> >> >> > array is passed to the outlined function
and stored in a
>> >> register inside
>> >> >> > the outlined function, but there's no
memory
allocation runtime
>> >> call.
>> >> >> > And in the device IR the array is passed
to the main
offloading
>> >> >> > function.
>> >> >> > Any suggestion or documentation on how
that memory
transfer is
>> >> >> > achieved?, and how to track its API
calls. It might be
that
>> I'm not
>> >> >> > using the proper omp directive, so here
is C code:
>> >> >>
>> >> >> Take a look at
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=yM9OaHxEHZBkU9HNFiMvQFv9zveVQp3kS%2FB4W9qkxKA%3D&reserved=0
>> >> , the mapping request,
>> >> >>
>> >> >> which may or may not cause data-transfers, are
part of the
call to
>> >> >> __tgt_target_teams[_nowait] and
>> __tgt_target_data_update[_nowait]. Start
>> >> >> by reading the impl. of those two functions in
the openmp
runtime
>> >> >> (llvm-project/openmp/...).
>> >> >>
>> >> >> Let me know if this helps and once you have new
questions.
>> >> >>
>> >> >> Cheers,
>> >> >> Johannes
>> >> >>
>> >> >> P.S. I merged your first patch. You didn't
run all test
before.
>> I did
>> >> >> and noticed that we failed one because of type
mismatches.
>> Fixed them
>> >> >> before I merged your patch.
>> >> >>
>> >> >> >
>> >> >> > #include <omp.h>
>> >> >> > #include <stdio.h>
>> >> >> >
>> >> >> > #define size 1000000
>> >> >> > double a[size];
>> >> >> >
>> >> >> > int main() {
>> >> >> >
>> >> >> > #pragma omp target map(to: a[0:size])
>> >> >> > #pragma omp teams distribute
>> >> >> > for (int i = 0; i < size; ++i) {
>> >> >> > a[i] = a[i]*(i*i)/2;
>> >> >> > }
>> >> >> >
>> >> >> > return 0;
>> >> >> > }
>> >> >> >
>> >> >> >
>> >> >> > PD: The command to compile is this:
>> >> >> > clang -v -save-temps \
>> >> >> > -I/path/to/llvm/release/9.x/include
-fopenmp \
>> >> >> > -fopenmp-targets=nvptx64-nvidia-cuda
-Xopenmp-target \
>> >> >> > -march=sm_61 memory_transfer.c -o
memory_transfer
>> >> >> >
>> >> >> > Thanks.
>> >> >> > Hamilton.
>> >> >> >
>> >> >> > On 23/03/20 11:11 a. m., Doerfert,
Johannes wrote:
>> >> >> >> On 3/22/20 12:33 PM, Hamilton Tobon
Mosquera wrote:
>> >> >> >> > Ahhh I understand. So I
believe the project you
posted
>> is still
>> >> >> >> opened. Do you any advice on the
question I posted
before?.
>> >> >> >>
>> >> >> >> Yes. Still open. You should use this
as the main part of
>> the GSoC
>> >> >> >> proposal (if you are interested). Can
you repeat the
question?
>> >> >> >>
>> >> >> >> >
>> >> >> >> > Thank you.
>> >> >> >> >
>> >> >> >> > On 22/03/20 1:26 p. m.,
Doerfert, Johannes wrote:
>> >> >> >> >> I did not. The issues
sections on GH was removed.
>> >> >> >> >>
>> >> >> >> >>
________________________________________
>> >> >> >> >> From: Hamilton Tobon
Mosquera <htobonm at eafit.edu.co>
>> >> >> >> >> Sent: Sunday, March 22,
2020 12:21
>> >> >> >> >> To: llvm-dev at
lists.llvm.org
>> >> >> >> >> Cc: Doerfert, Johannes
>> >> >> >> >> Subject: Re: [llvm-dev]
GSoC Interested student
>> >> >> >> >>
>> >> >> >> >> Hi there Johannes!,
>> >> >> >> >>
>> >> >> >> >> I was checking out the
issue (about latency
hiding when
>> >> >> transferring
>> >> >> >> >> memory from host to
device) you posted about a week
>> ago, but
>> >> >> it's not
>> >> >> >> >> there anymore. Did you
remove it?.
>> >> >> >> >>
>> >> >> >> >> Thanks.
>> >> >> >> >>
>> >> >> >> >> Hamilton.
>> >> >> >> >>
>> >> >> >> >> On 21/03/20 10:04 p. m.,
Hamilton Tobon Mosquera via
>> llvm-dev
>> >> >> wrote:
>> >> >> >> >>> Hi there!,
>> >> >> >> >>>
>> >> >> >> >>> Right now I'm
starting with the project proposed by
>> >> Johannes here:
>> >> >> >> >>>
>> >> >> >>
>> >> >>
>> >>
>>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.com%2Fllvm%2Fllvm-project%2Fissues%2F180&data=01%7C01%7Chtobonm%40eafit.edu.co%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=CVDD7HCzFVsPwSpZLOeiO5%2BMNX78sGrDd2eYJ5dWzyE%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=hN9WGEO%2F3quio7A9tvtIDHxznURuITYVBszFSmqM5M8%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=CVDD7HCzFVsPwSpZLOeiO5%2BMNX78sGrDd2eYJ5dWzyE%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=lycF%2FQ2WOJtIzRuGEgAu%2BOEw9220%2FnqslSkJMKFENJw%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gAMSEONp%2F8cMDuO0IckxHPXbBnUllJMWVfLHwa6UsBc%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=gAMSEONp%2F8cMDuO0IckxHPXbBnUllJMWVfLHwa6UsBc%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=Wye5YuFLVUPPESqTRnwsZtimjF0wdUxDetwoCNCib64%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=Wye5YuFLVUPPESqTRnwsZtimjF0wdUxDetwoCNCib64%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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%7C3e9ac7a8643c4dbe405e08d7d38dfd75%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=7CKfLhJkFKsjWn97SBMgKMz9nxRV1iZp305G9rsVLoA%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.
>> >
>>
>>
>> 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-Apr-25 14:49 UTC
[llvm-dev] GSoC Interested student
Hi Johannes,
I'm still writing the test cases with the FIXMEs that show how the
runtime call should be split and moved forward and backwards. I'm
starting with __tgt_target_teams and the possible situations that might
occur. The patch that Shilei did allows for making the computation
asynchronous too, not only the data transfer, it looks something like:
...
target_data_begin(..., __tgt_async_info a)
...
run_team_region(..., __tgt_async_info a)
target_data_end(..., __tgt_async_info a)
...
synchronize(__tgt_async_info a)
My question is, should the wait call contain target_data_run(...), that
is, running the target region only when calling wait()?, or should it
only call synchronize(...), that is, leaving the computation to the
issue(...) function?. The first case would look something like:
issue() {
...
target_data_begin()
...
}
wait() {
run_team_region() // The second approach has this function in the issue
target_data_end() // as well as this one.
synchronize()
}
The thing is that some behaviors change given the first or second
approach. For example:
heavyComputation() {
...
int a = rand();
...
int x = a % rand();
...
args[0] = &a;
issue(..., args, ...)
...
}
In that example, if the variable 'a' is shared and the 'issue'
function
has the 'run_team_region' call, it wouldn't be correct to move the
issue
past the expression 'int x = a % rand();', because a might be modified.
But if the issue only transfers data, there wouldn't be any problem I
guess. I think either approach has its advantages and disadvantages, for
example having the computation in the issue would prevent the transfer
to be moved even more forward. What do you think I should do?.
Thanks.
Hamilton.
On 7/04/20 11:25 p. m., Doerfert, Johannes wrote:> On 4/7/20 7:32 AM, Hamilton Tobon Mosquera wrote:
> > Hi Johannes,
> >
> > In order to create the testcases with the FIXMEs showing how it
> should look after the transformation, I've been trying to come up with
a
> way of splitting the runtime call __tgt_target_teams based on Silei's
> recent patch D77005. I like the idea of using its asynchronous version
> __tgt_target_teams_nowait, but I think we would have to modify it,
> either to return the async_info struct:
> >
> > __tgt_async_info __tgt_target_teams_nowait(...) // This would break
> the current canonical way of returning a status code.
> >
> > or to receive it as a pointer and modify its content, then waiting
on
> that async_info object:
> >
> > int __tgt_target_teams_nowait(..., __tgt_async_info *async_info >
nullptr)
> >
> > The first way leads to modify the Clang's code generation, and
I'm
> not sure if that is a good idea. Concerning the second way, I'm not
sure
> if having a default argument would lead also to modify the Clang's code
> generation.
> >
> > Another possible solution would be to implement another runtime
> function, avoiding the need to modify the _nowait version. This would
> lead to implement new runtime functions for each runtime call that
> implies data transfers, which is perhaps worse. I don't know which
> approach is better, should I look for a way of modifying Clang's code
> generation in case of needed?. I'm trying to be as less disruptive as
> possible. Can you give me some advice?.
>
> We are in the process of moving clangs OpenMP code generation into the
> OpenMPIRBuilder. While doing so we refactor/replace/rewrite stuff as we
> see fit. Thus, changing Clang, or better writing a different code
> generation in the OpenMPIRBuilder, is always an option.
>
> That said, I was expecting to create two new runtime calls per existing
> one but I have not really thought this through. What I imagined is
> something vaguely looking like this:
>
> Current:
> int omp_transfer_data(device, src, dst, size);
>
> New:
> ```
> // In the runtime {
> handle omp_transfer_data_issue(device, src, dst, size) {
> // some logic that was in omp_transfer_data before creating and using
> // handle
> return some_handle; // probably extended __tgt_async_info, can be an
> // argument as well.
> }
> int omp_transfer_data_wait(device, src, dst, size, handle) {
> // some logic that was in omp_transfer_data before using handle
> return error_code;
> }
> // }
>
> // In the runtime or in the header
> int omp_transfer_data(device, src, dst, size) {
> handle = omp_transfer_data_issue(...);
> return omp_transfer_data_wait(..., handle);
> }
> ```
>
> Now we can inline omp_transfer_data or replace it by the two calls to
> make the issue and wait explicit.
>
> Long story short, we will most certainly need new functions, how they
> look is up to you initially. Look at the other functions for inspiration
> and if something is off we'll let you know during code review ;)
>
> Does that help? Feel free to ask more questions!
>
> Cheers,
> Johannes
>
> P.S. Tomorrow there is an OpenMP in LLVM bi-weekly call which might be
> interesting (though it might be too early):
>
https://nam01.safelinks.protection.outlook.com/?url=http%3A%2F%2Flists.llvm.org%2Fpipermail%2Fopenmp-dev%2F2020-April%2F003292.html&data=01%7C01%7Chtobonm%40eafit.edu.co%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=ko3rPG8v3YMKHe4laYdIWfmOK%2Fl%2BSLjjtzTVldx8t%2FE%3D&reserved=0
>
>
> > Thank you.
> >
> > Hamilton.
> >
> > On 28/03/20 11:05 p. m., Doerfert, Johannes wrote:
> >> On 3/28/20 8:33 PM, Hamilton Tobon Mosquera wrote:
> >> > Hi Johannes,
> >> >
> >> > I've been thinking on how to implement a solution for
the
> problem and
> >> > I have some questions:
> >> >
> >> > 1. You mention in the problem description that the
current
> >> > asynchronous functions can be used, and probably
modified. The
> problem
> >> > is that those runtime functions with _nowait at the end
are not
> >> > asynchronous. Looking at their implementations, they just
wait for
> >> > other tasks is there are pending tasks, and then call its
> synchronous
> >> > versions. So, is it a good idea to wrap them in functions
that call
> >> > them as asynchronous (using std::async), then return an
object
> to wait
> >> > on?. Then, waiting on the object returned.
> >>
> >> Take a look at D77005. Shilei introduced some real askync
versions for a
> >> different reason already. You can build on top of that patch and
extend
> >> it if needed further.
> >>
> >>
> >> > 2. I'm thinking on replacing the data transfer inside
> >> > __tgt_target_teams with __tgt_target_data_begin. The
problem
> with this
> >> > is that __tgt_target_data_begin does not initialize
[first]private
> >> > variables, and doesn't check for lambda mappings.
Therefore, my idea
> >> > to go around this problem is to create another wrapper
function F,
> >> > that calls __tgt_target_data_begin, and executes the
necessary code
> >> > for lambda mappings check and [first]private variables,
and finally,
> >> > calling F as the asynchronous function. What is good
about this, is
> >> > that the code concerning data transfer inside
__tgt_target_teams can
> >> > be replaced with F.
> >>
> >> Interesting. I'll need to see small examples. For now, focus
on the
> >> simple case, we'll get to the rest later. However, we should
really
> >> start by creating test cases that show the initial code and
FIXMEs that
> >> describe how we want it to look like.
> >>
> >>
> >> > 3. Also, I was wondering if I must to be aware of (split
them into
> >> > issue and wait) all the other constructs that imply data
transfers
> >> > (update, data enter/exit), or just focus on
> __tgt_target[_nowait] and
> >> > __tgt_target_teams[_nowait].
> >>
> >> We want to eventually split all transfers into two but let's
start with
> >> a single one.
> >>
> >>
> >> > 4. This might be a dumb question, but just to make sure.
Do I
> have to
> >> > move the wait call into the offloaded region just before
the data is
> >> > used (because I think this would imply analyzing the
offloaded
> >> > bitcode, which is in another module)?, or move it just
before the
> >> > kernel is executed?.
> >>
> >> We split, then we move them apart. The cannot be moved over
instructions
> >> that might change behavior if we move them, e.g. an aliasing
memory
> >> operation or a kernel call that uses the data.
> >>
> >>
> >> > 5. Do you have a template for the GSoC proposal. That is,
is it
> plain
> >> > text or pdf or ...?. Also, do you have sort of a
structure that we
> >> > have to follow?, for example, Abstract -> Benefits
-> Deliverables
> >> > ..., or do we have freedom on that?. Your post in hte
GSoC page only
> >> > mentions that we must include our background, but nothing
about the
> >> > format of the proposal.
> >>
> >> I do not have one. I thought they might or they have an FAQ that
> >> describes what to do. I know they have some required things, so
check
> >> their webpage. I think you have to use google docs, I might be
wrong.
> >> Your categories look good, make the deliverables not too
specific but
> >> describe more ideas to work on. We'll make tasks as we go as
you are not
> >> the only one working on this general area.
> >>
> >> Feel free to send me a draft of the proposal for feedback.
> >>
> >> Cheers,
> >> Johannes
> >>
> >> > Thank you.
> >> >
> >> > Hamilton.
> >> >
> >> >
> >> > On``` 25/03/20 11:41 a. m., Doerfert, Johannes wrote:
> >> >> On 3/25/20 5:51 AM, Hamilton Tobon Mosquera wrote:
> >> >> > Hi Johannes,
> >> >> >
> >> >> > Sorry about the failing tests. I completely
forgot about
> running all
> >> >> > of them, I only ran the tests for the openmp
ipo transforms.
> >> >>
> >> >> No worries, it happens.
> >> >>
> >> >>
> >> >> > I found the GH issue in cache. I can post it
here if you want.
> >> >>
> >> >> Yes, please.
> >> >>
> >> >>
> >> >> > I analyzed these runtime calls,
> __tgt_target_teams_[nowait]. What I
> >> >> > don't understand about the problem we want
to solve
> (splitting the
> >> >> > memory transfer into issue and wait and
legally moving code
> >> inside) is
> >> >> > that the memory transfer is made inside those
runtime calls.
> >> >> > Therefore, the runtime calls to the memory
transfer cannot be
> >> analyzed
> >> >> > in the IR. Could you please explain a bit
further if the memory
> >> >> > transfer should split in the IR? or split
> >> __tgt_target_teams_[nowait]
> >> >> > in such a way that the memory transfer is made
in the first
> >> part, and
> >> >> > the wait in the second part.
> >> >>
> >> >> Exactly what you said in the last sentence. Split the
runtime
> call into
> >> >> two, issue and wait. Then we move issue
"up" the execution path
> and wait
> >> >> "down".
> >> >>
> >> >> Cheers,
> >> >> Johannes
> >> >>
> >> >>
> >> >> > Thanks.
> >> >> >
> >> >> > Hamilton.
> >> >> >
> >> >> > On 23/03/20 2:21 p. m., Doerfert, Johannes
wrote:
> >> >> >>
> >> >> >> On 3/23/20 11:22 AM, Hamilton Tobon
Mosquera wrote:
> >> >> >> > Hi Johannes,
> >> >> >> >
> >> >> >> > Yes, I'm highly interested, I
already started. It
> would be nice
> >> >> if you
> >> >> >> > post the details again. I
didn't take note of them, I was
> >> always
> >> >> >> > looking at the issue in GH.
> >> >> >>
> >> >> >> I asked if we can make the readonly but
unsure yet. I can
> write
> >> it again
> >> >> >> but it'll take ma some time.
> >> >> >>
> >> >> >>
> >> >> >> > The problem I'm facing right
now is this:
> >> >> >> > The first task is identifying
OpenMP API calls referred to
> >> >> device memory
> >> >> >> > allocation. I first understood how
the compiler driver
> >> >> orchestrates the
> >> >> >> > process of offloading OpenMP
regions to devices, and
> now I'm
> >> >> looking at
> >> >> >> > the step of the pipeline that
generates LLVM IR for both
> >> host and
> >> >> >> > device. I wrote a simple C program
that uses OpenMP to
> map a
> >> >> local array
> >> >> >> > to the device main memory. Then
I'm looking at the
> device and
> >> >> host LLVM
> >> >> >> > IR representation of this program,
but I can't find any
> >> runtime call
> >> >> >> > referred to device memory
allocation. In the host IR I see
> >> that the
> >> >> >> > array is passed to the outlined
function and stored in a
> >> >> register inside
> >> >> >> > the outlined function, but
there's no memory
> allocation runtime
> >> >> call.
> >> >> >> > And in the device IR the array is
passed to the main
> offloading
> >> >> >> > function.
> >> >> >> > Any suggestion or documentation on
how that memory
> transfer is
> >> >> >> > achieved?, and how to track its API
calls. It might be
> that
> >> I'm not
> >> >> >> > using the proper omp directive, so
here is C code:
> >> >> >>
> >> >> >> Take a look at
> >> >>
> >>
>
https://nam01.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgodbolt.org%2Fz%2FE9VoVx&data=01%7C01%7Chtobonm%40eafit.edu.co%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=OrktJWgp%2Bl9BL6fTBZAeiEOfLZ8cD8IYy7OuiLcq3A8%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pgWPly2cBQJ7YlJPsKAk4KvFqdZ7rrXYFdCTwlYeT3U%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=ZRXzDI%2FMTOYsGd4%2BRIljHCrWTECjLlz%2F8zvO7TolDMU%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pgWPly2cBQJ7YlJPsKAk4KvFqdZ7rrXYFdCTwlYeT3U%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=Gx6hd7a5%2FAxDAP%2BJVSnkSGujoWWWPQ2BO%2Fw0Ch53L5c%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=E0VJfvDVQ3euU3%2BVTNrFWSdgkGEwBGQeOEohIumeBRU%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=E0VJfvDVQ3euU3%2BVTNrFWSdgkGEwBGQeOEohIumeBRU%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=agKF%2FzZ7EHND5JiuiMV5W84wH9R1mStsoJxanpAeQXg%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=agKF%2FzZ7EHND5JiuiMV5W84wH9R1mStsoJxanpAeQXg%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pBCrDFOV3a917meFVEW3IgmdQumSohVtBPBGly32TXI%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pBCrDFOV3a917meFVEW3IgmdQumSohVtBPBGly32TXI%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pBCrDFOV3a917meFVEW3IgmdQumSohVtBPBGly32TXI%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%7Ca0688315910b434282a408d7db6c8c8a%7C99f7b55e9cbe467b8143919782918afb%7C0&sdata=pBCrDFOV3a917meFVEW3IgmdQumSohVtBPBGly32TXI%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.
> >> >
> >>
> >>
> >> 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.