May 2018 - Anl-intel-lrz-knr - mailman3.cels.anl.gov

Re: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug!
by Tannenbaum, Barry M 31 May '18

31 May '18

Please post it. From: anl-intel-csa [mailto:[email protected]] On Behalf Of Applencourt, Thomas Sent: Thursday, May 31, 2018 4:19 PM To: Velesko, Paulius <pvelesko(a)anl.gov> Cc: anl-intel-csa(a)lists.jlse.anl.gov Subject: Re: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug! Hi Paulius, I have indeed access to this forum. I will post my detailed bug report there if nobody has any objection. Regards, Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036 ________________________________ From: Velesko, Paulius <paulius.velesko(a)intel.com<mailto:[email protected]>> Sent: Thursday, May 31, 2018 14:56 To: Applencourt, Thomas; anl-intel-csa(a)lists.jlse.anl.gov<mailto:[email protected]> Subject: RE: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug! Hi Argonne, According to the wiki page you have 2 options Do you have access to the forum? (ask barry.m.tannenbaum(a)intel.com<mailto:[email protected]> for access) or you can post it here https://wiki.jlse.anl.gov/display/intelcsa/Bug%2c+Issue+and+Request+Tracker Best, Paulius From: anl-intel-csa [mailto:[email protected]] On Behalf Of Applencourt, Thomas Sent: Thursday, May 31, 2018 2:48 PM To: anl-intel-csa(a)lists.jlse.anl.gov<mailto:[email protected]> Subject: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug! Hi Intel, I found a probable bug in the Functional simulator. For security reasons, I cannot join any files. What is the procedure to bring this information up to you? I may be able to send an AES-256 encrypted zip containing all the data needed to reproduce the bug to the people who need it (please correct me if I'm wrong). I can follow any other procedures you think are more relevant, Regards, Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036

1 0

Re: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug!
by Applencourt, Thomas 31 May '18

31 May '18

Hi Paulius, I have indeed access to this forum. I will post my detailed bug report there if nobody has any objection. Regards, Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036 ________________________________ From: Velesko, Paulius <paulius.velesko(a)intel.com> Sent: Thursday, May 31, 2018 14:56 To: Applencourt, Thomas; anl-intel-csa(a)lists.jlse.anl.gov Subject: RE: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug! Hi Argonne, According to the wiki page you have 2 options Do you have access to the forum? (ask barry.m.tannenbaum(a)intel.com for access) or you can post it here https://wiki.jlse.anl.gov/display/intelcsa/Bug%2c+Issue+and+Request+Tracker Best, Paulius From: anl-intel-csa [mailto:[email protected]] On Behalf Of Applencourt, Thomas Sent: Thursday, May 31, 2018 2:48 PM To: anl-intel-csa(a)lists.jlse.anl.gov Subject: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug! Hi Intel, I found a probable bug in the Functional simulator. For security reasons, I cannot join any files. What is the procedure to bring this information up to you? I may be able to send an AES-256 encrypted zip containing all the data needed to reproduce the bug to the people who need it (please correct me if I'm wrong). I can follow any other procedures you think are more relevant, Regards, Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036

1 0

Re: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug!
by Velesko, Paulius 31 May '18

31 May '18

Hi Argonne, According to the wiki page you have 2 options Do you have access to the forum? (ask barry.m.tannenbaum(a)intel.com for access) or you can post it here https://wiki.jlse.anl.gov/display/intelcsa/Bug%2c+Issue+and+Request+Tracker Best, Paulius From: anl-intel-csa [mailto:[email protected]] On Behalf Of Applencourt, Thomas Sent: Thursday, May 31, 2018 2:48 PM To: anl-intel-csa(a)lists.jlse.anl.gov Subject: [Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug! Hi Intel, I found a probable bug in the Functional simulator. For security reasons, I cannot join any files. What is the procedure to bring this information up to you? I may be able to send an AES-256 encrypted zip containing all the data needed to reproduce the bug to the people who need it (please correct me if I'm wrong). I can follow any other procedures you think are more relevant, Regards, Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036

1 0

[Anl-intel-csa] > csa fusion pass: Internal error in fusion code - probably a bug!
by Applencourt, Thomas 31 May '18

31 May '18

Hi Intel, I found a probable bug in the Functional simulator. For security reasons, I cannot join any files. What is the procedure to bring this information up to you? I may be able to send an AES-256 encrypted zip containing all the data needed to reproduce the bug to the people who need it (please correct me if I'm wrong). I can follow any other procedures you think are more relevant, Regards, Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036

1 0

[Anl-intel-csa] Canceled: ANL/Intel CSA Workload Characterization
by Petrini, Fabrizio 30 May '18

30 May '18

Placeholder for weekly ANL/Intel CSA Exascale workload characterization meeting ......................................................................................................................................... --> Join Skype Meeting<https://meet.intel.com/fabrizio.petrini/7K7FM32F> Trouble Joining? Try Skype Web App<https://meet.intel.com/fabrizio.petrini/7K7FM32F?sl=1> Join by phone +1(916)356-2663 (or your local bridge access #) Choose bridge 5.<tel:+1(916)356-2663%20(or%20your%20local%20bridge%20access%20#)%20Choose%20bridge%205.> (Global) English (United States) Find a local number<https://dial.intel.com> Conference ID: 791574018 Forgot your dial-in PIN?<https://dial.intel.com> |Help<https://o15.officeredir.microsoft.com/r/rlidLync15?clid=1033&p1=5&p2=2009> [!OC([1033])!] .........................................................................................................................................

1 0

[Anl-intel-csa] Our conference call this Friday at 1:00pm central
by Kumaran, Kalyan 30 May '18

30 May '18

I’m canceling this week as I’m on vacation and it does not look like Intel is ready to present on Xeons. -Kumar

1 0

Re: [Anl-intel-csa] Please explain cache miss result
by Bertoni, Colleen 30 May '18

30 May '18

Hi Zheming, Thanks for the interesting example! It was a good example to try to understand what was happening. While I was working through it, I found something that might be useful about the map clause for pointers. For the n=1 case, if we use “#pragma omp target map(to: a) map(from: c)” as in the example, there are 4 memory accesses for the n=1 case (0 hits, 4 misses, like you reported), since it looks like a and c are passed by reference. (Verified by looking at the graph and .loads file.) When the map clause is changed to “#pragma omp target map(to: a[0:n]) map(from: c[0:n])”, there are fewer memory accesses (0 hits, 2 misses for n=1), since it seems like the addresses of a and c are just passed as addresses. (Verified by looking at the graph and .loads file.) I think this means that we should always specify the size of the pointer if we’re using it as an array. I think the OMP 4.0 standard document says to do this anyway, but I don’t remember seeing this happen in the past for CSA code, although I could be wrong. (The current CSA-Getting-Started-Gude.pdf specifies “In this release, when compiling for CSA target, every pointer variable that is used in the lexical scope of a "target" construct should be listed in a "map" clause. There is no need to specify the size of the pointee.” However, here it seems like we need to specify the size of the pointee to avoid the extra memory accesses. )  Thanks, Colleen ________________________________ From: anl-intel-csa <anl-intel-csa-bounces(a)lists.jlse.anl.gov> on behalf of Jin, Zheming <zjin(a)anl.gov> Sent: Tuesday, May 29, 2018 11:38 AM To: anl-intel-csa(a)jlse.anl.gov Subject: [Anl-intel-csa] Please explain cache miss result I was motivated by Colleen's cache analysis in the simple memory model. However, I don't understand the cache miss results from the .stat file for the following program. Thanks for the light. void example (int n, const int * __restrict__ a, int * __restrict__ c) { int i; #pragma omp target map(to: a) map(from: c) { for (i=0; i<n; i++) { c[i] = a[i] + 1; } } 14 hits, 10 misses, n=11 15 hits, 7 misses, n= 10 13 hits, 7 misses, n=9 11 hits, 7 misses, n=8 9 hits, 7 misses, n=7 5 hits, 7 misses, n=5 3 hits, 7 misses, n=4 2 hits, 6 misses, n=3 1 hits, 5 misses, n=2 0 hits, 4 misses, n=1

1 0

[Anl-intel-csa] CSA Architecture Reference Typo
by Applencourt, Thomas 30 May '18

30 May '18

Hi Intel, I found one typo in the CSA Architecture Reference Rev 0.8-6: - page 92 in the sample usage of the `nestrepeat1` unit, the `pick1` should have been a `switch1`. More precisely, replace the line: ``` pick1 %ign, innerincr, outeri, 1 ```` by ``` switch1 %ign, innerincr, outeri, 1 ``` or arguably better: ``` filter1 innerincr, outeri ``` Thanks for this document, is a mine of information! Thomas Applencourt Leadership Computing Facility Argonne National Laboratory Building 240 Office - 1.D.18 9700 S Cass Av., Argonne Il, 60439 (630) 252-0036

1 0

Re: [Anl-intel-csa] Atomic Operations and Visibility Outputs
by Hal Finkel 30 May '18

30 May '18

On 04/23/2018 06:36 PM, Hal Finkel wrote: > Hi, Simon, Kent, Elliott, et al., > > In Santa Clara we had a brief discussion about the lowering of C/C++ > atomic operations to the CSA. I'll try to summarize that discussion and > highlight one significant concern... > > In C++ there are several memory ordering choices for atomic operations > (C is similar): > > 1. Relaxed - These impose no extra ordering requirements and the > local-visibility output is all that is required (i.e., these can be > treated just like non-atomic memory accesses). > > 2. Consume - We didn't discuss this in detail, but it is generally > considered unimplementable as specified outside of mapping it onto > Acquire (so I'll assume we'll do so). > > 3. Acquire - No reads or writes in the current thread can be reordered > before this load. For the CSA, this means that we'll use the > global-visibility output instead of the local-visibility output for all > successors. > > 3. Release - No reads or writes in the current thread can be reordered > after this store. For the CSA, this means that the store needs to depend > on the global visibility of all predecessors, or alternatively, must be > preceded by a fence. > > 4. Acquire/Release - For an operations that both reads and writes. The > requirements are combined. > > 5. Sequentially Consistent - All such atomic operations also form a > total order,. For CSA, this means that each seq. cst. atomic operation > must depend on the global visibility of all seq. cst. predecessors, or > alternatively, must be preceded by a fence. > > In the current CSA model, Relaxed and Acquire (and Consume) are trivial > to implement. Additional sequential consistency semantics also seem easy > (it just adds another dependency chain). Release is problematic, > however. A release store needs to depend on the global visibility of all > predecessors (to be clear, this means all accesses including the > non-atomic ones). This can either be done directly (i.e., by forming a > separate sub-DAG of global-dependency tokens to use and the dependency > input to the release store) or implicitly by just preceding each such > release store with a fence operation. > > The fence operation is easy, but expensive, and Kent said his desire is > to do something similar to what I've done in our prototype compiler > (which is somehow track the global-visibility outputs of preceding > memory operations) for finer-grained control (and, presumably, > significantly lower overhead). While this might, at first, seem > equivalent to the fence, it's not, because we might have multiple > independent workers on the CSA and a fence in any worker will affect all > workers (whereas separate tracking won't have that problem). Also, in a > pipelined situation, a fence in the pipeline will force > logically-earlier iterations in the pipeline to wait on the completion > of memory accesses from logically-later iterations in the pipeline > (which are now in flight by the time that the earlier iterations hit the > fence at the bottom of the pipeline). > > Now the problem: Simon told us that, at least for CSA v1, each memory > operation will likely have only one visibility output. You won't be able > to get both the local and global visibility output from a memory > operation. This seems to make the finer-grained dependency tracking much > less appealing (and, perhaps, even slower than using the fence). The > problem is, if we do that, no separate tracking is possible (because we > get only one visibility output, we'd need to always use the global one, > and we'd need to use that as the input for all subsequent operations > (not just the store)). This is going to make it impossible to hide any > of the latency of the regular memory accesses, even far away from the > atomic store, and all of the memory accesses will be very slow as a > result. This seems pretty awful. > > If this only-one-visibility-output restriction could be lifted, that > would seem much better. As it stands, with this restriction, I don't see > a path to high-performance code using atomic operations (e.g., lock-free > data structures or critical sections) on the CSA. Any thoughts on this? Thanks again, Hal > > Thanks again, > > Hal > -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory

1 0

[Anl-intel-csa] converting omp_get_thread_num() automatically
by Jin, Zheming 29 May '18

29 May '18

I think it will be useful for the ICX compiler to support automatic conversion of the following histogram program in the future release. Users may still need to reset the 2D array, local_bins, to zeroes outside the OpenMP target region. 15 #pragma omp parallel 16 { 17 int id = omp_get_thread_num(); 18 int i, j; 19 20 #pragma omp for 21 for(i = 0 ; i < n ; ++i) 22 local_bins[id][data[i] / k] ++; 23 24 #pragma omp for 25 for(i = 0 ; i < num_bins; ++i) 26 { 27 for(j = 0 ; j < number_of_threads; ++j) 28 { 29 bins[i] += local_bins[j][i]; 30 } 31 } 32 }

1 0