similar to: [RFC] Late (OpenMP) GPU code "SPMD-zation"

<font size=2 face="sans-serif">Hi Johannes,</font><br><br><font size=2 face="sans-serif">Thank you for the explanation.</font><br><br><font size=2 face="sans-serif">I think we need to clarify some details about code generation in Clang today:</font><br><br><font size=2

We are working on OpenMP target offloading for GPUs in Flang, and adopting the same code generation strategy. The proposal is affecting us. It would be nice to know more details about the proposal. So we can prepare ourselves to adapt flang (if everything goes on the way). Have you find and a solution for data sharing? How are you going to manage data sharing for SPMD and non-SPMD? From: cfe-dev

We would still know that. We can do exactly the same reasoning as we do now. I think the important question is, how different is the code generated for either mode and can we hide (most of) the differences in the runtime. If I understand you correctly, you say the data sharing code looks very different and the differences cannot be hidden, correct? It would be helpful for me to understand your

We could still do that in clang, couldn't we? Get Outlook for Android<https://aka.ms/ghei36> ________________________________ From: Alexey Bataev <a.bataev at outlook.com> Sent: Tuesday, January 22, 2019 12:52:42 PM To: Doerfert, Johannes Rudolf; cfe-dev at lists.llvm.org Cc: openmp-dev at lists.llvm.org; LLVM-Dev; Finkel, Hal J.; Alexey Bataev; Arpith Chacko Jacob Subject: Re:

1. You don't need to implement everything in a single patch. The development process is a step-by-step process, when you commit something in small pieces. The code must nit be fully functional, you may start from some basic features. Currently it is very hard to review. 2. I rather doubt that it can be reused without changes for AMD etc., especially without being fully tested. The only tested

Could you elaborate on what you refer to wrt data sharing. What do we currently do in the clang code generation that we could not effectively implement in the runtime, potentially with support of an llvm pass. Thanks, James Get Outlook for Android<https://aka.ms/ghei36> ________________________________ From: Alexey Bataev <a.bataev at outlook.com> Sent: Tuesday, January 22, 2019

Johannes, did you try it on AMD GPUs? If not, I think it might be early to claim it as a general interface for NVidia/AMD GPUs. I'm ok, if you want tointroduce a basic class for the GPU-specific codegen, but it must be done step-by-step and thoroughly tested and reviewed. Theremightbe some parts, common with NVPTX codegen. You can put the commonfunctions into a base class and remove them from

There are tooooooo(!) many changes, I don't who's going to review sooooo big patch. You definitely need to split it into several smaller patches. Also, I don't like the idea adding of one more class for NVPTX codegen. All your changes should be on top of the eixisting solution. ------------- Best regards, Alexey Bataev 13.03.2019 15:08, Doerfert, Johannes пишет: > Please consider

------------- Best regards, Alexey Bataev 13.03.2019 15:35, Doerfert, Johannes пишет: > > Hi Alexey, > > > thank you for your quick feedback. > > > > There are tooooooo(!) many changes, I don't who's going to review sooooo big > patch. > > > I can for sure split it in the three components/repositories that are > touched, clang, llvm, and openmp.

Yes, I understand. I just don't understand why there are 15 review links. as to the format of the list, we previously had a list of feature for 4.5 (before it was fully supported), we can reuse this old format. Best regards, Alexey Bataev > 6 июня 2019 г., в 19:12, Narayanaswamy, Ravi <ravi.narayanaswamy at intel.com> написал(а): > > Alexey, > Johannes want to put all

I don't know where it comes from, currently there is just 2 patches with the new functionality : mapper implementation in clang and in the runtime. Plus the patch for unified memory. Best regards, Alexey Bataev > 6 июня 2019 г., в 19:03, Doerfert, Johannes <jdoerfert at anl.gov> написал(а): > >> On 06/06, Alexey Bataev via llvm-dev wrote: >> Hmm, it is interesting.

Hmm, it is interesting. What's missing? and why it is "a lot". It is the status for clang 8.0. There were not too many changes in trunk since last release. Best regards, Alexey Bataev 6 июня 2019 г., в 8:24, Doerfert, Johannes via llvm-dev <llvm-dev at lists.llvm.org<mailto:llvm-dev at lists.llvm.org>> написал(а): Thanks for pointing that one out to me! It's a

Long story short: We want to create a webpage under llvm.org that shows the implementation status of OpenMP features, similar to https://clang.llvm.org/cxx_status.html . Please let me know if you have wishes, questions, or concerns. Cheers, Johannes

For example, Type 2 case, scalar-foo used call by value while vector-foo used call by ref. The question Johannes is asking is whether we can decipher that after the fact, only by looking at the two function signatures, or need some more info (what kind, what's minimal)? I think we need to list up cases of interest, and for each vector ABI of interest, we need to work on the requirements and

Hi, [I started this discussion on a single mailing list (flang-dev [0]) but this is a notice to the others (llvm, clang, openmp) so we hopefully get all interested parties involved.] This is an RFC for the design of the OpenMP front-ends under the LLVM umbrella. It is necessary to talk about this now as Flang (aka. F18) is maturing at a very promising rate and about to become a sub-project

> On Jun 24, 2019, at 10:53 AM, Tian, Xinmin <xinmin.tian at intel.com> wrote: > > To me, it is also an issue related to SIMD signature matching when the vectorizer kicks in. Losing info from FE to BE is not good in general. > Yes, we cannot loose such information. In particular, the three examples I reported are all generating i64 in the scalar function signature: // Type 1

@Xinmin, Saito: If Clang/the frontend generates the version there is no problem, or is there? The frontend knows about the original source type and it's ABI specific lowering already. @Francesco, we should even consider putting the generating capabilities outside of the OpenMP code generation (in the future). That could allow easier reuse by other frontends. Get Outlook for

Hi All, I have a doubt regarding the behavior of LoopAccessAnalysis on incorrect #pragma omp simd with -fopenmp-simd flag. How should the compiler behave if the #pragma omp simd on a loop is incorrect and can be proved by Loop Access Analysis. Here is the sample code. #pragma omp simd for (dim_t p = 0; p < m; ++p) #pragma unroll for (dim_t i = 0; i < 6; ++i) { {

>In all cases, the IR type of the parameters in `foo` is i64, therefore is not possible to distinguish what C type generated the signature of `foo`. Ouch. >I don’t know if this is going to be a problem for other architectures I haven't checked what IA-32/Intel64 should do for type 2, but I fully agree that this needs to be done properly according to the ABI. >Therefore, I would

>Thank you everybody for their input, and for your patience. This is proving harder than expected! :) Thank you for doing the hard part of the work. Hideki -----Original Message----- From: Francesco Petrogalli [mailto:Francesco.Petrogalli at arm.com] Sent: Monday, June 24, 2019 11:26 AM To: Saito, Hideki <hideki.saito at intel.com> Cc: Doerfert, Johannes <jdoerfert at anl.gov>;