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

[Anl-intel-csa] Update on fsim not stoping after the out_ctl lic is populated
by Applencourt, Thomas 28 Feb '18

28 Feb '18

Hi, In a previous email (see issue #1 https://wiki.jlse.anl.gov/display/intelcsa/Bug%2C+Issue+and+Request+Tracker here), I report a case where the functional simulator using the relaxed mode was not exiting after showing the message: csasim: warning: one or more LICs not empty on exit. I was able to reproduce this behavior in the following minimal example: .text .globl queue .entry queue, hybriddataflow .set implicit .set relaxed .result .lic .nil out_ctl .param .lic .nil in_ctl .param .lic .i32 n repeat1 bit, 1, n # will produce and print indefinitely debugprint bit copy1 out_ctl, %ign, %ign, %ign, in_ctl Thank you, Regards, Thomas 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] redaddfN
by Hammond, Jeff R 28 Feb '18

28 Feb '18

Intel Fortran and C/C++ compilers re-associate parentheses (i.e. the -funsafe-math-optimizations part of -ffast-math) at every level of optimization unless the user specifically disables this (with -fp-model=source or -assume protect_parens) so any DOE codes that use the Intel compilers are at least somewhat insensitive to floating-point summation order. I am surprised that most Argonne applications are not using -ffast-math or equivalent. This can be worth as much as 1.3x in total application performance for codes that make heavy use of division and/or transcendentals. Molecular dynamics codes like LAMMPS enable this option by default. It's probably the easiest double-digit performance speedup one can get without code changes, other than linking against a vendor BLAS library… If numerical accuracy of reductions is a concern, you may want to investigate the use of Kahan summation, which is implementable at near-zero overhead in CSA. There's an implementation in the CSA examples directory already. Jeff From: anl-intel-csa <anl-intel-csa-bounces(a)lists.jlse.anl.gov> on behalf of Hal Finkel <hfinkel(a)anl.gov> Organization: Argonne National Laboratory Date: Wednesday, February 28, 2018 at 10:51 AM To: "Kermin (Elliott) Fleming" <kermin.fleming(a)intel.com>, "Kumaran, Kalyan" <kumaran(a)anl.gov>, "anl-intel-csa(a)jlse.anl.gov" <anl-intel-csa(a)jlse.anl.gov> Subject: Re: [Anl-intel-csa] redaddfN Hi, Elliot, An obvious concern with software mitigation using individual fadds is that, unless there's some trick to keeping the effective ordering the same, the mitigation can't be applied automatically if the application is not compiled with -ffast-math (or similar) or uses special pragmas on the loops in question. Many of our current applications are not compiled this way. Thanks again, Hal On 02/28/2018 11:11 AM, Fleming, Kermin wrote: Unfortunately, this is the expected behavior due to implementation choices in the FMAs. The upside of FMAs is that they are more dense and have some energy/op improvement. The downside is that the multiply and add pipelines get mixed together. The single-cycle reduction we had aimed for previously was a result of separating out multiply and add functionality. We will continue to look in to the provisioning of an FAdd PE, which could support fast reduction. This will depend on software’s ability to cope. I’m sure Kent will comment on software mitigation here, but the main thrust is loop pipelining. -Elliott From: anl-intel-csa [mailto:[email protected]] On Behalf Of Kumaran, Kalyan Sent: Wednesday, February 28, 2018 10:49 AM To: anl-intel-csa(a)jlse.anl.gov<mailto:[email protected]> Subject: [Anl-intel-csa] Fwd: redaddfN Kent, Elliott, could one of you explain this and if this is the expected behavior in real hardware: o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. Thanks, Kumar Begin forwarded message: From: Vitali A Morozov <morozov(a)anl.gov<mailto:[email protected]>> Subject: Fwd: Re: redaddfN Date: February 27, 2018 at 11:16:13 AM CST To: "Kumaran, Kalyan" <kumaran(a)alcf.anl.gov<mailto:[email protected]>> Kumar, FYI. A short version of this thread: in a recent Feb 22 update, Intel has changed the behavior of a fp reduction unit. It is not pipelined anymore, accepting input once every 4 cycles. Most of us think this change will have significant impact on performance and should be inverted back. Thomas has implemented a temporal workaround to replace this unit with a pipelined graph at the expense of complexity and latency. Vitali -------- Forwarded Message -------- Subject: Re: redaddfN Date: Tue, 27 Feb 2018 10:45:30 -0600 From: Applencourt, Thomas <tapplencourt(a)anl.gov><mailto:[email protected]> To: Morozov, Vitali A. <morozov(a)alcf.anl.gov><mailto:[email protected]> Yes. You can find the implementation here: `/soft/restricted/intel_csa/csatools/csa-p/lib/libreduc.S` The implementation is straightforward. I use four readaddf32 in a round-robin fashion and do a manual reduction add the end. The only `trick`, is to duplicate the end signal to be able to stop all the 4 reduction unit. For that, I use the `replace` unit. To be totally honest, this is not a strict equivalent of the `readdf32` unit because I reorder the operation during the reduction, so numerical instability may occur. It's the price to pay to be able to pipeline. 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: Morozov, Vitali A. Sent: Tuesday, February 27, 2018 10:38 To: Applencourt, Thomas Subject: Re: redaddfN Hi Thomas, if I understand this thread correctly, you have implemented a pipelined reduction graph with the functionality, similar to readaddf32 unit. Is this correct? Vitali On 2/23/18 11:33 AM, Hal Finkel wrote: Hi, Thomas, Can you share the code? -Hal On 02/22/2018 04:32 PM, Applencourt, Thomas wrote: Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.)§ All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected This will have a big impact on application reeling on redaddf32N. So I did my own implementation [1] of this unit who allow pipelining like the old one. This new unit is a drop-in replacement of the old one. You can find the timing bellow: old : 4476 new_intel: 16678 new_anl : 4493 Regards, [1]: Hal, I dit it. I use one of their `strange` unit, the replace1 one. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 11:57 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version The discovery of new units continue. · completion (completion{8,16,32,64}availidx.Cd.u8, ldres.Cd.iN, stidx.Cu.u8, stdata.Cu.iN, size.Lu.u8) ``` Another mode of operation for scratchpads are as a “completion buffer”, where values can be inserted in a “sliding window”, and values can be retrieved in order. The completion operation supports this functionality. Note that this is a very unusual operator in terms of how it behaves. It can be viewed as effectively 3 independent operations bound together. I'm not yet understanding this unit. I have the feeling that it's a important one. - Snull (snull res.CRd.i1, ctlseq.CRLu.il<http://ctlseq.crlu.il/>) determine if current stream has zero length. Needed I think to implemented a more efficient `reduction add`. - Filter (filter{0-64} res.CRd.iN, filter.CRu.i1, val.CRLu.iN) This operation filters a set of values based on corresponding predicates. If filter is true, val is copied to res. If filter is false, val is consumed, but res is not written. syntactic sugar for switchN %ign, r, f, v - They also add a flag in the {pick/switch}any unit. This flag is used to change the mode of selection (priority, or round-robin). pickany{0-64} dst.CRd.iN, idx.CRd.i1, i0.CRLu.iN, i1.CRLu.iN, mode.Lu.i1=0 If the mode operand has the value 0 (default), this performs a prioritized pick, selecting i0 preferentially over i1 if both are available. If the mode operand has value 1, this will pick the last operand not selected, if available. ``` Implementing a Load balance strategy is now far more easy. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 09:19 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version Note that they add the documentation for 2 additional units: · prefetch (yeah!) · guard (sync a value with a signal) They also document new MEMLEVEL values (used by the memory ops): MEMLEVEL_T0 Use all levels of the cache hierarchy as appropriate MEMLEVEL_T1 Avoid any very local cache, if present, but use the main CSA cache. For CSA Version 1, there is no difference from using MEMLEVEL_T0 MEMLEVEL_T2 Bypass the first two levels of cache. For CSA Version 1, this implies avoiding the tile level cache. MEMLEVEL_NTA This asserts there is no benefit to caching (Non-Temporal Always), and references should bypass the (note: architectural cache hierarchy to the degree possible. NOTE: This is not just a hint, and things like memory ordering operands may not provide the same guarantees when this is used. 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: anl-intel-csa anl-intel-csa-bounces(a)lists.jlse.anl.gov<http://mailto:[email protected]/> on behalf of Glossop, Kent kent.glossop(a)intel.com<http://mailto:[email protected]/> Intel | Data Center Solutions, IoT, and PC Innovation<http://mailto:kent%[email protected]/> intel.com<http://intel.com/> Intel's innovation in cloud computing, data center, Internet of Things, and PC solutions is powering the smart and connected digital world we live in. Sent: Wednesday, February 21, 2018 22:54 To: Glossop, Kent Subject: [Anl-intel-csa] CSA SDK Release Notes for 2018/02/21 version CSA SDK Release Notes for 2018/02/21 version Intel Corporation Proprietary The default release has been changed to 2018-02-21-release. Relevant changes since the 2017-01-22-release: Documentation · The CSA Architecture document has finally been updated in the doc. area to reflect new operations and update operation signatures. It should be current with the simulator in this release. We expect to keep them synchronized in the future. The first entry in the edit history after the table of contents has a list of the additions/changes, and there are change bars in the margin. Spatial Advisor This includes a major update of Spatial Advisor which now incorporates per-application and per-loopnest offload strategy analyzer, L2 cache and HBM simple performance modeling and numerous other updates. New features: · Introduced analysis across overall program and all loopnests, not just innermost loops · Stabilized loopnest offload strategy selector · Per-program (Amdhal’s law) area and speed-up metrics introduced · L2 + HBM (external memory) performance model added; cache simulation enabled. · Performance estimate consideration of offload costs, such as configuration, invocation, and (optional) data transfer costs · New outputs, including a graphical representation of the call tree with offloaded, non-accelerable and offloadable regions highlighted · More precise area model for “Heavy INT” operations · Introduced “bound by” CSA workloads characterization to estimate fractions of PE-, memory-, cache- bound code. Compiler / Runtime · Inner loop Pipelining o Inner loop pipelining code generation now specifies minimum backedge LIC depths on the output assembly. This avoids deadlock even if the simulator creates shallow LICs by default. (Previously, ILPL assumed that all LICs would have a depth of 128.) o Inner loop pipelining code generation now waits for all of a gang of cohorts to be available before admitting them. This eliminates the possibility of deadlock in certain cases where buffering outside of the inner loop is insufficient. However, in these cases, performance will not be optimal due to the insufficient buffering. Future work on buffer insertion should address this. · SPMDization o Improved SPMDization intrinsic detection, resolving some crashes and incorrect code generation in cases involving fully unrolled loops. o Added a warning for loops where blocked SPMDization was requested but where the trip count cannot be determined by the compiler, which use to cause crashes. · Memory reference ordering o Fixed ordering of sequential outer loops containing parallel inner loops. o Fixed the handling of instructions that both read and write memory, which were previously mistakenly ordered as loads. o Re-implemented dropping of redundant dependencies from a particular memory instruction back to itself. o Removed old memory ordering passes and some of the old internal options for accessing them. At one point, we had recommended passing “-mllvm -csa-order-memops-type=wavefront” (one of the options that has been removed by this change) in order to avoid suboptimal handling of SPMDization by what was then the latest memory ordering pass. Since the current pass does not have that problem, this is no longer necessary and any uses of that option can be dropped. · Atomics o The compiler now emits atomics as “compare and swap” loops with atmcmpxchg rather than single “read modify write” operations such as “atmadd”. This better aligns with current hardware expectations. For experimentation purposes, the RMW-style operations can be restored by adding -Xclang -target-feature -Xclang ‘+rmwatomic’ to CFLAGS/CPPFLAGS/CSA_COMPILE_OPTIONS. Note that the simulator timing of atomic operations is still considered optimistic. o Operands to atmcmpxchg instructions are now emitted in the correct order Simulator · New Features o Addition of snull, stpick, ststream operations o The assembler now verifies that all scopes are properly closed o Stats added to track memory accesses that are not naturally aligned · Modified Features o Stream comparison operators changed to also return a control sequence: cmprelT cmpres.CRd.i1, valres.CRd.T, ctlseqres.CRd.i1, ctlseqa.CRLu.i1, vala.CRLu.T, ctlseqb.CRLu.i1, valb.CRLu.T · Removed Features - a number of deprecated operations have been removed from the simulator o Signed scalar loads (ldsN, ldsNd, ldsNx) have been removed. The unsigned scalar loads (ldN, ldNd, ldNx) are unchanged o All operations that used f34 or f66 datatypes have been removed · Pending Removal o The old “wide” memory references (ldNx8/stNx8) are deprecated and will be going away in the future. These should not be confused with the wide streaming memory operations (sldNx8/sstNx8), which are effectively superset replacements. o The seqc operation will be removed from the examples and the simulator as well. Code that uses them should be converted. o While they will not be removed, the atomics other than atmswp and atmcmpswp are not expected to be supported for V1 and should not be used for code. If you have a compelling example where these would make a significant performance difference to applications, please bring that to our attention. · Fixed Bugs o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected o Atmcmpxchg now returns the correct value · Known Issues o Atomic operations are modeled as a memory read o The fence operation is currently not invoking the memory model to wait for memory o The simple memory model is not waiting until memory quiesces before returning from a function. -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory _______________________________________________ anl-intel-csa mailing list anl-intel-csa(a)lists.jlse.anl.gov<mailto:[email protected]> https://lists.jlse.anl.gov/mailman/listinfo/anl-intel-csa -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory

1 0

Re: [Anl-intel-csa] redaddfN
by Fleming, Kermin 28 Feb '18

28 Feb '18

A purpose-built FAdd supporting reduction would maintain order. The problem is design complexity, but we will keep pushing on a definition here. The existing proposal of software loop pipelining may not maintain order in all cases. It is helpful to quantify your pessimism on software mitigation. How bad is forcing -ffast-math? That sounds like a coarse and dangerous approach. Of course, annotation also sounds painful. -Elliott From: Hal Finkel [mailto:[email protected]] Sent: Wednesday, February 28, 2018 1:51 PM To: Fleming, Kermin <kermin.fleming(a)intel.com>; Kumaran, Kalyan <kumaran(a)anl.gov>; anl-intel-csa(a)jlse.anl.gov Subject: Re: [Anl-intel-csa] redaddfN Hi, Elliot, An obvious concern with software mitigation using individual fadds is that, unless there's some trick to keeping the effective ordering the same, the mitigation can't be applied automatically if the application is not compiled with -ffast-math (or similar) or uses special pragmas on the loops in question. Many of our current applications are not compiled this way. Thanks again, Hal On 02/28/2018 11:11 AM, Fleming, Kermin wrote: Unfortunately, this is the expected behavior due to implementation choices in the FMAs. The upside of FMAs is that they are more dense and have some energy/op improvement. The downside is that the multiply and add pipelines get mixed together. The single-cycle reduction we had aimed for previously was a result of separating out multiply and add functionality. We will continue to look in to the provisioning of an FAdd PE, which could support fast reduction. This will depend on software’s ability to cope. I’m sure Kent will comment on software mitigation here, but the main thrust is loop pipelining. -Elliott From: anl-intel-csa [mailto:[email protected]] On Behalf Of Kumaran, Kalyan Sent: Wednesday, February 28, 2018 10:49 AM To: anl-intel-csa(a)jlse.anl.gov<mailto:[email protected]> Subject: [Anl-intel-csa] Fwd: redaddfN Kent, Elliott, could one of you explain this and if this is the expected behavior in real hardware: o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. Thanks, Kumar Begin forwarded message: From: Vitali A Morozov <morozov(a)anl.gov<mailto:[email protected]>> Subject: Fwd: Re: redaddfN Date: February 27, 2018 at 11:16:13 AM CST To: "Kumaran, Kalyan" <kumaran(a)alcf.anl.gov<mailto:[email protected]>> Kumar, FYI. A short version of this thread: in a recent Feb 22 update, Intel has changed the behavior of a fp reduction unit. It is not pipelined anymore, accepting input once every 4 cycles. Most of us think this change will have significant impact on performance and should be inverted back. Thomas has implemented a temporal workaround to replace this unit with a pipelined graph at the expense of complexity and latency. Vitali -------- Forwarded Message -------- Subject: Re: redaddfN Date: Tue, 27 Feb 2018 10:45:30 -0600 From: Applencourt, Thomas <tapplencourt(a)anl.gov><mailto:[email protected]> To: Morozov, Vitali A. <morozov(a)alcf.anl.gov><mailto:[email protected]> Yes. You can find the implementation here: `/soft/restricted/intel_csa/csatools/csa-p/lib/libreduc.S` The implementation is straightforward. I use four readaddf32 in a round-robin fashion and do a manual reduction add the end. The only `trick`, is to duplicate the end signal to be able to stop all the 4 reduction unit. For that, I use the `replace` unit. To be totally honest, this is not a strict equivalent of the `readdf32` unit because I reorder the operation during the reduction, so numerical instability may occur. It's the price to pay to be able to pipeline. 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: Morozov, Vitali A. Sent: Tuesday, February 27, 2018 10:38 To: Applencourt, Thomas Subject: Re: redaddfN Hi Thomas, if I understand this thread correctly, you have implemented a pipelined reduction graph with the functionality, similar to readaddf32 unit. Is this correct? Vitali On 2/23/18 11:33 AM, Hal Finkel wrote: Hi, Thomas, Can you share the code? -Hal On 02/22/2018 04:32 PM, Applencourt, Thomas wrote: Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.)§ All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected This will have a big impact on application reeling on redaddf32N. So I did my own implementation [1] of this unit who allow pipelining like the old one. This new unit is a drop-in replacement of the old one. You can find the timing bellow: old : 4476 new_intel: 16678 new_anl : 4493 Regards, [1]: Hal, I dit it. I use one of their `strange` unit, the replace1 one. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 11:57 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version The discovery of new units continue. · completion (completion{8,16,32,64}availidx.Cd.u8, ldres.Cd.iN, stidx.Cu.u8, stdata.Cu.iN, size.Lu.u8) ``` Another mode of operation for scratchpads are as a “completion buffer”, where values can be inserted in a “sliding window”, and values can be retrieved in order. The completion operation supports this functionality. Note that this is a very unusual operator in terms of how it behaves. It can be viewed as effectively 3 independent operations bound together. I'm not yet understanding this unit. I have the feeling that it's a important one. - Snull (snull res.CRd.i1, ctlseq.CRLu.il<http://ctlseq.crlu.il/>) determine if current stream has zero length. Needed I think to implemented a more efficient `reduction add`. - Filter (filter{0-64} res.CRd.iN, filter.CRu.i1, val.CRLu.iN) This operation filters a set of values based on corresponding predicates. If filter is true, val is copied to res. If filter is false, val is consumed, but res is not written. syntactic sugar for switchN %ign, r, f, v - They also add a flag in the {pick/switch}any unit. This flag is used to change the mode of selection (priority, or round-robin). pickany{0-64} dst.CRd.iN, idx.CRd.i1, i0.CRLu.iN, i1.CRLu.iN, mode.Lu.i1=0 If the mode operand has the value 0 (default), this performs a prioritized pick, selecting i0 preferentially over i1 if both are available. If the mode operand has value 1, this will pick the last operand not selected, if available. ``` Implementing a Load balance strategy is now far more easy. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 09:19 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version Note that they add the documentation for 2 additional units: · prefetch (yeah!) · guard (sync a value with a signal) They also document new MEMLEVEL values (used by the memory ops): MEMLEVEL_T0 Use all levels of the cache hierarchy as appropriate MEMLEVEL_T1 Avoid any very local cache, if present, but use the main CSA cache. For CSA Version 1, there is no difference from using MEMLEVEL_T0 MEMLEVEL_T2 Bypass the first two levels of cache. For CSA Version 1, this implies avoiding the tile level cache. MEMLEVEL_NTA This asserts there is no benefit to caching (Non-Temporal Always), and references should bypass the (note: architectural cache hierarchy to the degree possible. NOTE: This is not just a hint, and things like memory ordering operands may not provide the same guarantees when this is used. 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: anl-intel-csa anl-intel-csa-bounces(a)lists.jlse.anl.gov<http://mailto:[email protected]/> on behalf of Glossop, Kent kent.glossop(a)intel.com<http://mailto:[email protected]/> Intel | Data Center Solutions, IoT, and PC Innovation<http://mailto:kent%[email protected]/> intel.com<http://intel.com/> Intel's innovation in cloud computing, data center, Internet of Things, and PC solutions is powering the smart and connected digital world we live in. Sent: Wednesday, February 21, 2018 22:54 To: Glossop, Kent Subject: [Anl-intel-csa] CSA SDK Release Notes for 2018/02/21 version CSA SDK Release Notes for 2018/02/21 version Intel Corporation Proprietary The default release has been changed to 2018-02-21-release. Relevant changes since the 2017-01-22-release: Documentation · The CSA Architecture document has finally been updated in the doc. area to reflect new operations and update operation signatures. It should be current with the simulator in this release. We expect to keep them synchronized in the future. The first entry in the edit history after the table of contents has a list of the additions/changes, and there are change bars in the margin. Spatial Advisor This includes a major update of Spatial Advisor which now incorporates per-application and per-loopnest offload strategy analyzer, L2 cache and HBM simple performance modeling and numerous other updates. New features: · Introduced analysis across overall program and all loopnests, not just innermost loops · Stabilized loopnest offload strategy selector · Per-program (Amdhal’s law) area and speed-up metrics introduced · L2 + HBM (external memory) performance model added; cache simulation enabled. · Performance estimate consideration of offload costs, such as configuration, invocation, and (optional) data transfer costs · New outputs, including a graphical representation of the call tree with offloaded, non-accelerable and offloadable regions highlighted · More precise area model for “Heavy INT” operations · Introduced “bound by” CSA workloads characterization to estimate fractions of PE-, memory-, cache- bound code. Compiler / Runtime · Inner loop Pipelining o Inner loop pipelining code generation now specifies minimum backedge LIC depths on the output assembly. This avoids deadlock even if the simulator creates shallow LICs by default. (Previously, ILPL assumed that all LICs would have a depth of 128.) o Inner loop pipelining code generation now waits for all of a gang of cohorts to be available before admitting them. This eliminates the possibility of deadlock in certain cases where buffering outside of the inner loop is insufficient. However, in these cases, performance will not be optimal due to the insufficient buffering. Future work on buffer insertion should address this. · SPMDization o Improved SPMDization intrinsic detection, resolving some crashes and incorrect code generation in cases involving fully unrolled loops. o Added a warning for loops where blocked SPMDization was requested but where the trip count cannot be determined by the compiler, which use to cause crashes. · Memory reference ordering o Fixed ordering of sequential outer loops containing parallel inner loops. o Fixed the handling of instructions that both read and write memory, which were previously mistakenly ordered as loads. o Re-implemented dropping of redundant dependencies from a particular memory instruction back to itself. o Removed old memory ordering passes and some of the old internal options for accessing them. At one point, we had recommended passing “-mllvm -csa-order-memops-type=wavefront” (one of the options that has been removed by this change) in order to avoid suboptimal handling of SPMDization by what was then the latest memory ordering pass. Since the current pass does not have that problem, this is no longer necessary and any uses of that option can be dropped. · Atomics o The compiler now emits atomics as “compare and swap” loops with atmcmpxchg rather than single “read modify write” operations such as “atmadd”. This better aligns with current hardware expectations. For experimentation purposes, the RMW-style operations can be restored by adding -Xclang -target-feature -Xclang ‘+rmwatomic’ to CFLAGS/CPPFLAGS/CSA_COMPILE_OPTIONS. Note that the simulator timing of atomic operations is still considered optimistic. o Operands to atmcmpxchg instructions are now emitted in the correct order Simulator · New Features o Addition of snull, stpick, ststream operations o The assembler now verifies that all scopes are properly closed o Stats added to track memory accesses that are not naturally aligned · Modified Features o Stream comparison operators changed to also return a control sequence: cmprelT cmpres.CRd.i1, valres.CRd.T, ctlseqres.CRd.i1, ctlseqa.CRLu.i1, vala.CRLu.T, ctlseqb.CRLu.i1, valb.CRLu.T · Removed Features - a number of deprecated operations have been removed from the simulator o Signed scalar loads (ldsN, ldsNd, ldsNx) have been removed. The unsigned scalar loads (ldN, ldNd, ldNx) are unchanged o All operations that used f34 or f66 datatypes have been removed · Pending Removal o The old “wide” memory references (ldNx8/stNx8) are deprecated and will be going away in the future. These should not be confused with the wide streaming memory operations (sldNx8/sstNx8), which are effectively superset replacements. o The seqc operation will be removed from the examples and the simulator as well. Code that uses them should be converted. o While they will not be removed, the atomics other than atmswp and atmcmpswp are not expected to be supported for V1 and should not be used for code. If you have a compelling example where these would make a significant performance difference to applications, please bring that to our attention. · Fixed Bugs o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected o Atmcmpxchg now returns the correct value · Known Issues o Atomic operations are modeled as a memory read o The fence operation is currently not invoking the memory model to wait for memory o The simple memory model is not waiting until memory quiesces before returning from a function. -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory _______________________________________________ anl-intel-csa mailing list anl-intel-csa(a)lists.jlse.anl.gov<mailto:[email protected]> https://lists.jlse.anl.gov/mailman/listinfo/anl-intel-csa -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory

1 0

Re: [Anl-intel-csa] redaddfN
by Hal Finkel 28 Feb '18

28 Feb '18

Hi, Elliot, An obvious concern with software mitigation using individual fadds is that, unless there's some trick to keeping the effective ordering the same, the mitigation can't be applied automatically if the application is not compiled with -ffast-math (or similar) or uses special pragmas on the loops in question. Many of our current applications are not compiled this way. Thanks again, Hal On 02/28/2018 11:11 AM, Fleming, Kermin wrote: > > Unfortunately, this is the expected behavior due to implementation > choices in the FMAs. The upside of FMAs is that they are more dense > and have some energy/op improvement. The downside is that the > multiply and add pipelines get mixed together. The single-cycle > reduction we had aimed for previously was a result of separating out > multiply and add functionality. We will continue to look in to the > provisioning of an FAdd PE, which could support fast reduction. This > will depend on software’s ability to cope. > > > > I’m sure Kent will comment on software mitigation here, but the main > thrust is loop pipelining. > > > > -Elliott > > > > *From:*anl-intel-csa [mailto:[email protected]] > *On Behalf Of *Kumaran, Kalyan > *Sent:* Wednesday, February 28, 2018 10:49 AM > *To:* anl-intel-csa(a)jlse.anl.gov > *Subject:* [Anl-intel-csa] Fwd: redaddfN > > > > Kent, Elliott, could one of you explain this and if this is the > expected behavior in real hardware: > > > > o Floating point reducer latency has changed, which may > require recoding of reductions to avoid a significant > performance hit. (Often it will be desirable to pipeline > multiple partial sums through a normal operation that can > accept a new value per cycle, and reduce the partial sums at > the end, rather than using a reduction operator.) > > § All floating point reducers now have a latency of 4 cycles, > and only accept a value once every 4 cycles. > > > > Thanks, > > Kumar > > > > Begin forwarded message: > > > > *From: *Vitali A Morozov <morozov(a)anl.gov <mailto:[email protected]>> > > *Subject: Fwd: Re: redaddfN* > > *Date: *February 27, 2018 at 11:16:13 AM CST > > *To: *"Kumaran, Kalyan" <kumaran(a)alcf.anl.gov > <mailto:[email protected]>> > > > > Kumar, > > > > FYI. A short version of this thread: in a recent Feb 22 update, > Intel has changed the behavior of a fp reduction unit. It is not > pipelined anymore, accepting input once every 4 cycles. > > > > Most of us think this change will have significant impact on > performance and should be inverted back. Thomas has implemented a > temporal workaround to replace this unit with a pipelined graph at > the expense of complexity and latency. > > > > Vitali > > > > -------- Forwarded Message -------- > > *Subject:* > > > > Re: redaddfN > > *Date:* > > > > Tue, 27 Feb 2018 10:45:30 -0600 > > *From:* > > > > Applencourt, Thomas <tapplencourt(a)anl.gov> > <mailto:[email protected]> > > *To:* > > > > Morozov, Vitali A. <morozov(a)alcf.anl.gov> > <mailto:[email protected]> > > > > Yes. > > > > You can find the implementation here: > `/soft/restricted/intel_csa/csatools/csa-p/lib/libreduc.S` > > > > The implementation is straightforward. I use four readaddf32 in a > round-robin fashion and do a manual reduction add the end. > > The only `trick`, is to duplicate the end signal to be able to > stop all the 4 reduction unit. For that, I use the `replace` unit. > > > > To be totally honest, this is not a strict equivalent of the > `readdf32` unit because I reorder the operation during the > reduction, so numerical instability may occur. It's the price to > pay to be able to pipeline. > > > > > > 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:* Morozov, Vitali A. > *Sent:* Tuesday, February 27, 2018 10:38 > *To:* Applencourt, Thomas > *Subject:* Re: redaddfN > > > > Hi Thomas, > > > > if I understand this thread correctly, you have implemented a > pipelined reduction graph with the functionality, similar to > readaddf32 unit. Is this correct? > > > > Vitali > > > > > > > > On 2/23/18 11:33 AM, Hal Finkel wrote: > > Hi, Thomas, > > > > Can you share the code? > > > > -Hal > > > > On 02/22/2018 04:32 PM, Applencourt, Thomas wrote: > > Floating point reducer latency has changed, which > may require recoding of reductions to avoid a > significant performance hit. (Often it will be > desirable to pipeline multiple partial sums > through a normal operation that can accept a new > value per cycle, and reduce the partial sums at > the end, rather than using a reduction operator.)§ > All floating point reducers now have a latency of > 4 cycles, and only accept a value once every 4 cycles. > § redaddfN, redsubfN, sredaddfN, sredsubfN latency > was erroneously 1 cycle. It is now 4 cycles like > the rest of the floating point reducers > § The declaration of external output operand for > memory operations has been corrected > > This will have a big impact on application reeling > on |redaddf32N|. So I did my own implementation [1] of > this unit who allow pipelining like the old one. This new > unit is a drop-in replacement of the old one. > > You can find the timing bellow: > > |old : 4476| > > |new_intel: 16678| > > |new_anl : 4493| > > Regards, > > [1]: Hal, I dit it. I use one of their `strange` unit, > the |replace1| one. > > 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: Applencourt, Thomas > Sent: Thursday, February 22, 2018 11:57 > To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. > Cc: Bertoni, Colleen; Jin, Zheming > Subject: Re: CSA SDK Release Notes for 2018/02/21 version > > The discovery of new units continue. > > · completion (completion{8,16,32,64}availidx.Cd.u8, > ldres.Cd.iN, stidx.Cu.u8, stdata.Cu.iN, size.Lu.u8) > > ``` > Another mode of operation for scratchpads are as a > “completion buffer”, where values can be inserted in a > “sliding window”, and values can be retrieved in order. > > The completion operation supports this functionality. Note > that this is a very unusual operator in terms of how it > behaves. It can be viewed as effectively 3 independent > operations bound together. > > |I'm not yet understanding this unit. I have the feeling > that it's a important one.| > > | | > > | | > > |- Snull (snull res.CRd.i1, |ctlseq.CRLu.il > <http://ctlseq.crlu.il/>|)| > > |determine if current stream has zero length. | > > | Needed I think to implemented a more efficient `reduction > add`.| > > | | > > | | > > |- Filter (filter{0-64} res.CRd.iN, filter.CRu.i1, > val.CRLu.iN)| > > |This operation filters a set of values based on > corresponding predicates. If filter is true, val is copied > to res. If filter is false, val is consumed, but res is > not written. | > > | syntactic sugar for switchN %ign, r, f, v | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > |- They also add a flag in the {pick/switch}any unit. This > flag is used to change the mode of selection (priority, or > round-robin).| > > |pickany{0-64} dst.CRd.iN, idx.CRd.i1, i0.CRLu.iN, > i1.CRLu.iN, mode.Lu.i1=0 | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > | | > > |If the mode operand has the value 0 (default), this > performs a prioritized pick, selecting i0 preferentially > over i1 if both are available. If the mode operand has > value 1, this will pick the last operand not selected, if > available. | > > ``` > > Implementing a Load balance strategy is now far more easy. > > 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: Applencourt, Thomas > Sent: Thursday, February 22, 2018 09:19 > To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. > Cc: Bertoni, Colleen; Jin, Zheming > Subject: Re: CSA SDK Release Notes for 2018/02/21 version > > Note that they add the documentation for 2 additional units: > > · prefetch (yeah!) > > · guard (sync a value with a signal) > > They also document new |MEMLEVEL| values (used by the > memory ops): > > MEMLEVEL_T0 > Use all levels of the cache hierarchy as appropriate > MEMLEVEL_T1 > Avoid any very local cache, if present, but use the main > CSA cache. For CSA Version 1, there is no > difference from using MEMLEVEL_T0 > MEMLEVEL_T2 > Bypass the first two levels of cache. For CSA Version 1, > this implies avoiding the tile level cache. > MEMLEVEL_NTA > This asserts there is no benefit to caching (Non-Temporal > Always), and references should bypass the > (note: architectural cache hierarchy to the degree > possible. NOTE: This is not just a hint, and things like > memory > ordering operands may not provide the same guarantees when > this is used. > > 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: > anl-intel-csa anl-intel-csa-bounces(a)lists.jlse.anl.gov > <http://mailto:[email protected]/> on > behalf of Glossop, Kent kent.glossop(a)intel.com > <http://mailto:[email protected]/> > > Intel | Data Center Solutions, IoT, and PC Innovation > <http://mailto:kent%[email protected]/> > > intel.com <http://intel.com/> > > Intel's innovation in cloud computing, data center, > Internet of Things, and PC solutions is powering the smart > and connected digital world we live in. > > > Sent: Wednesday, February 21, 2018 22:54 > To: Glossop, Kent > Subject: [Anl-intel-csa] CSA SDK Release Notes for > 2018/02/21 version > > CSA SDK Release Notes for 2018/02/21 version > > Intel Corporation Proprietary > > The default release has been changed to 2018-02-21-release. > > Relevant changes since the 2017-01-22-release: > > Documentation > > · The CSA Architecture document has finally been > updated in the doc. area to reflect new operations and > update operation signatures. It should be current with the > simulator in this release. We expect to keep them > synchronized in the future. The first entry in the edit > history after the table of contents has a list of the > additions/changes, and there are change bars in the margin. > > Spatial Advisor > > This includes a major update of Spatial Advisor which now > incorporates per-application and per-loopnest offload > strategy analyzer, L2 cache and HBM simple performance > modeling and numerous other updates. New features: > > · Introduced analysis across overall program and > all loopnests, not just innermost loops > > · Stabilized loopnest offload strategy selector > > · Per-program (Amdhal’s law) area and speed-up > metrics introduced > > · L2 + HBM (external memory) performance model > added; cache simulation enabled. > > · Performance estimate consideration of offload > costs, such as configuration, invocation, and (optional) > data transfer costs > > · New outputs, including a graphical representation > of the call tree with offloaded, non-accelerable and > offloadable regions highlighted > > · More precise area model for “Heavy INT” operations > > · Introduced “bound by” CSA workloads > characterization to estimate fractions of PE-, memory-, > cache- bound code. > > Compiler / Runtime > > · Inner loop Pipelining > > o Inner loop pipelining code generation now specifies > minimum backedge LIC depths on the output assembly. This > avoids deadlock even if the simulator creates shallow LICs > by default. (Previously, ILPL assumed that all LICs would > have a depth of 128.) > > o Inner loop pipelining code generation now waits for all > of a gang of cohorts to be available before admitting > them. This eliminates the possibility of deadlock in > certain cases where buffering outside of the inner loop is > insufficient. However, in these cases, performance will > not be optimal due to the insufficient buffering. Future > work on buffer insertion should address this. > > · SPMDization > > o Improved SPMDization intrinsic detection, resolving some > crashes and incorrect code generation in cases involving > fully unrolled loops. > > o Added a warning for loops where blocked SPMDization was > requested but where the trip count cannot be determined by > the compiler, which use to cause crashes. > > · Memory reference ordering > > o Fixed ordering of sequential outer loops containing > parallel inner loops. > > o Fixed the handling of instructions that both read and > write memory, which were previously mistakenly ordered as > loads. > > o Re-implemented dropping of redundant dependencies from a > particular memory instruction back to itself. > > o Removed old memory ordering passes and some of the old > internal options for accessing them. At one point, we had > recommended passing “-mllvm > -csa-order-memops-type=wavefront” (one of the options that > has been removed by this change) in order to avoid > suboptimal handling of SPMDization by what was then the > latest memory ordering pass. Since the current pass does > not have that problem, this is no longer necessary and any > uses of that option can be dropped. > > · Atomics > > o The compiler now emits atomics as “compare and swap” > loops with atmcmpxchg rather than single “read modify > write” operations such as “atmadd”. This better aligns > with current hardware expectations. For experimentation > purposes, the RMW-style operations can be restored by > adding -Xclang -target-feature -Xclang ‘+rmwatomic’ to > CFLAGS/CPPFLAGS/CSA_COMPILE_OPTIONS. Note that the > simulator timing of atomic operations is still considered > optimistic. > > o Operands to atmcmpxchg instructions are now emitted in > the correct order > > Simulator > > · New Features > > o Addition of snull, stpick, ststream operations > > o The assembler now verifies that all scopes are properly > closed > > o Stats added to track memory accesses that are not > naturally aligned > > · Modified Features > > o Stream comparison operators changed to also return a > control sequence: > > cmprelT cmpres.CRd.i1, valres.CRd.T, ctlseqres.CRd.i1, > ctlseqa.CRLu.i1, vala.CRLu.T, ctlseqb.CRLu.i1, valb.CRLu.T > > · Removed Features - a number of deprecated > operations have been removed from the simulator > > o Signed scalar loads (ldsN, ldsNd, ldsNx) have been > removed. The unsigned scalar loads (ldN, ldNd, ldNx) are > unchanged > > o All operations that used f34 or f66 datatypes have been > removed > > · Pending Removal > > o The old “wide” memory references (ldNx8/stNx8) are > deprecated and will be going away in the future. These > should not be confused with the wide streaming memory > operations (sldNx8/sstNx8), which are effectively superset > replacements. > > o The seqc operation will be removed from the examples and > the simulator as well. Code that uses them should be > converted. > > o While they will not be removed, the atomics other than > atmswp and atmcmpswp are not expected to be supported for > V1 and should not be used for code. If you have a > compelling example where these would make a significant > performance difference to applications, please bring that > to our attention. > > · Fixed Bugs > > o Floating point reducer latency has changed, which may > require recoding of reductions to avoid a significant > performance hit. (Often it will be desirable to pipeline > multiple partial sums through a normal operation that can > accept a new value per cycle, and reduce the partial sums > at the end, rather than using a reduction operator.) > > § All floating point reducers now have a latency of 4 > cycles, and only accept a value once every 4 cycles. > > § redaddfN, redsubfN, sredaddfN, sredsubfN latency was > erroneously 1 cycle. It is now 4 cycles like the rest of > the floating point reducers > > § The declaration of external output operand for memory > operations has been corrected > > o Atmcmpxchg now returns the correct value > > · Known Issues > > o Atomic operations are modeled as a memory read > > o The fence operation is currently not invoking the memory > model to wait for memory > > o The simple memory model is not waiting until memory > quiesces before returning from a function. > > > > -- > > Hal Finkel > > Lead, Compiler Technology and Programming Languages > > Leadership Computing Facility > > Argonne National Laboratory > > > > > > > > > > _______________________________________________ > anl-intel-csa mailing list > anl-intel-csa(a)lists.jlse.anl.gov > https://lists.jlse.anl.gov/mailman/listinfo/anl-intel-csa -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory

1 0

Re: [Anl-intel-csa] redaddfN
by Fleming, Kermin 28 Feb '18

28 Feb '18

Unfortunately, this is the expected behavior due to implementation choices in the FMAs. The upside of FMAs is that they are more dense and have some energy/op improvement. The downside is that the multiply and add pipelines get mixed together. The single-cycle reduction we had aimed for previously was a result of separating out multiply and add functionality. We will continue to look in to the provisioning of an FAdd PE, which could support fast reduction. This will depend on software’s ability to cope. I’m sure Kent will comment on software mitigation here, but the main thrust is loop pipelining. -Elliott From: anl-intel-csa [mailto:[email protected]] On Behalf Of Kumaran, Kalyan Sent: Wednesday, February 28, 2018 10:49 AM To: anl-intel-csa(a)jlse.anl.gov Subject: [Anl-intel-csa] Fwd: redaddfN Kent, Elliott, could one of you explain this and if this is the expected behavior in real hardware: o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. Thanks, Kumar Begin forwarded message: From: Vitali A Morozov <morozov(a)anl.gov<mailto:[email protected]>> Subject: Fwd: Re: redaddfN Date: February 27, 2018 at 11:16:13 AM CST To: "Kumaran, Kalyan" <kumaran(a)alcf.anl.gov<mailto:[email protected]>> Kumar, FYI. A short version of this thread: in a recent Feb 22 update, Intel has changed the behavior of a fp reduction unit. It is not pipelined anymore, accepting input once every 4 cycles. Most of us think this change will have significant impact on performance and should be inverted back. Thomas has implemented a temporal workaround to replace this unit with a pipelined graph at the expense of complexity and latency. Vitali -------- Forwarded Message -------- Subject: Re: redaddfN Date: Tue, 27 Feb 2018 10:45:30 -0600 From: Applencourt, Thomas <tapplencourt(a)anl.gov><mailto:[email protected]> To: Morozov, Vitali A. <morozov(a)alcf.anl.gov><mailto:[email protected]> Yes. You can find the implementation here: `/soft/restricted/intel_csa/csatools/csa-p/lib/libreduc.S` The implementation is straightforward. I use four readaddf32 in a round-robin fashion and do a manual reduction add the end. The only `trick`, is to duplicate the end signal to be able to stop all the 4 reduction unit. For that, I use the `replace` unit. To be totally honest, this is not a strict equivalent of the `readdf32` unit because I reorder the operation during the reduction, so numerical instability may occur. It's the price to pay to be able to pipeline. 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: Morozov, Vitali A. Sent: Tuesday, February 27, 2018 10:38 To: Applencourt, Thomas Subject: Re: redaddfN Hi Thomas, if I understand this thread correctly, you have implemented a pipelined reduction graph with the functionality, similar to readaddf32 unit. Is this correct? Vitali On 2/23/18 11:33 AM, Hal Finkel wrote: Hi, Thomas, Can you share the code? -Hal On 02/22/2018 04:32 PM, Applencourt, Thomas wrote: Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.)§ All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected This will have a big impact on application reeling on redaddf32N. So I did my own implementation [1] of this unit who allow pipelining like the old one. This new unit is a drop-in replacement of the old one. You can find the timing bellow: old : 4476 new_intel: 16678 new_anl : 4493 Regards, [1]: Hal, I dit it. I use one of their `strange` unit, the replace1 one. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 11:57 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version The discovery of new units continue. · completion (completion{8,16,32,64}availidx.Cd.u8, ldres.Cd.iN, stidx.Cu.u8, stdata.Cu.iN, size.Lu.u8) ``` Another mode of operation for scratchpads are as a “completion buffer”, where values can be inserted in a “sliding window”, and values can be retrieved in order. The completion operation supports this functionality. Note that this is a very unusual operator in terms of how it behaves. It can be viewed as effectively 3 independent operations bound together. I'm not yet understanding this unit. I have the feeling that it's a important one. - Snull (snull res.CRd.i1, ctlseq.CRLu.il<http://ctlseq.crlu.il/>) determine if current stream has zero length. Needed I think to implemented a more efficient `reduction add`. - Filter (filter{0-64} res.CRd.iN, filter.CRu.i1, val.CRLu.iN) This operation filters a set of values based on corresponding predicates. If filter is true, val is copied to res. If filter is false, val is consumed, but res is not written. syntactic sugar for switchN %ign, r, f, v - They also add a flag in the {pick/switch}any unit. This flag is used to change the mode of selection (priority, or round-robin). pickany{0-64} dst.CRd.iN, idx.CRd.i1, i0.CRLu.iN, i1.CRLu.iN, mode.Lu.i1=0 If the mode operand has the value 0 (default), this performs a prioritized pick, selecting i0 preferentially over i1 if both are available. If the mode operand has value 1, this will pick the last operand not selected, if available. ``` Implementing a Load balance strategy is now far more easy. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 09:19 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version Note that they add the documentation for 2 additional units: · prefetch (yeah!) · guard (sync a value with a signal) They also document new MEMLEVEL values (used by the memory ops): MEMLEVEL_T0 Use all levels of the cache hierarchy as appropriate MEMLEVEL_T1 Avoid any very local cache, if present, but use the main CSA cache. For CSA Version 1, there is no difference from using MEMLEVEL_T0 MEMLEVEL_T2 Bypass the first two levels of cache. For CSA Version 1, this implies avoiding the tile level cache. MEMLEVEL_NTA This asserts there is no benefit to caching (Non-Temporal Always), and references should bypass the (note: architectural cache hierarchy to the degree possible. NOTE: This is not just a hint, and things like memory ordering operands may not provide the same guarantees when this is used. 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: anl-intel-csa anl-intel-csa-bounces(a)lists.jlse.anl.gov<http://mailto:[email protected]/> on behalf of Glossop, Kent kent.glossop(a)intel.com<http://mailto:[email protected]/> Intel | Data Center Solutions, IoT, and PC Innovation<http://mailto:kent%[email protected]/> intel.com<http://intel.com/> Intel's innovation in cloud computing, data center, Internet of Things, and PC solutions is powering the smart and connected digital world we live in. Sent: Wednesday, February 21, 2018 22:54 To: Glossop, Kent Subject: [Anl-intel-csa] CSA SDK Release Notes for 2018/02/21 version CSA SDK Release Notes for 2018/02/21 version Intel Corporation Proprietary The default release has been changed to 2018-02-21-release. Relevant changes since the 2017-01-22-release: Documentation · The CSA Architecture document has finally been updated in the doc. area to reflect new operations and update operation signatures. It should be current with the simulator in this release. We expect to keep them synchronized in the future. The first entry in the edit history after the table of contents has a list of the additions/changes, and there are change bars in the margin. Spatial Advisor This includes a major update of Spatial Advisor which now incorporates per-application and per-loopnest offload strategy analyzer, L2 cache and HBM simple performance modeling and numerous other updates. New features: · Introduced analysis across overall program and all loopnests, not just innermost loops · Stabilized loopnest offload strategy selector · Per-program (Amdhal’s law) area and speed-up metrics introduced · L2 + HBM (external memory) performance model added; cache simulation enabled. · Performance estimate consideration of offload costs, such as configuration, invocation, and (optional) data transfer costs · New outputs, including a graphical representation of the call tree with offloaded, non-accelerable and offloadable regions highlighted · More precise area model for “Heavy INT” operations · Introduced “bound by” CSA workloads characterization to estimate fractions of PE-, memory-, cache- bound code. Compiler / Runtime · Inner loop Pipelining o Inner loop pipelining code generation now specifies minimum backedge LIC depths on the output assembly. This avoids deadlock even if the simulator creates shallow LICs by default. (Previously, ILPL assumed that all LICs would have a depth of 128.) o Inner loop pipelining code generation now waits for all of a gang of cohorts to be available before admitting them. This eliminates the possibility of deadlock in certain cases where buffering outside of the inner loop is insufficient. However, in these cases, performance will not be optimal due to the insufficient buffering. Future work on buffer insertion should address this. · SPMDization o Improved SPMDization intrinsic detection, resolving some crashes and incorrect code generation in cases involving fully unrolled loops. o Added a warning for loops where blocked SPMDization was requested but where the trip count cannot be determined by the compiler, which use to cause crashes. · Memory reference ordering o Fixed ordering of sequential outer loops containing parallel inner loops. o Fixed the handling of instructions that both read and write memory, which were previously mistakenly ordered as loads. o Re-implemented dropping of redundant dependencies from a particular memory instruction back to itself. o Removed old memory ordering passes and some of the old internal options for accessing them. At one point, we had recommended passing “-mllvm -csa-order-memops-type=wavefront” (one of the options that has been removed by this change) in order to avoid suboptimal handling of SPMDization by what was then the latest memory ordering pass. Since the current pass does not have that problem, this is no longer necessary and any uses of that option can be dropped. · Atomics o The compiler now emits atomics as “compare and swap” loops with atmcmpxchg rather than single “read modify write” operations such as “atmadd”. This better aligns with current hardware expectations. For experimentation purposes, the RMW-style operations can be restored by adding -Xclang -target-feature -Xclang ‘+rmwatomic’ to CFLAGS/CPPFLAGS/CSA_COMPILE_OPTIONS. Note that the simulator timing of atomic operations is still considered optimistic. o Operands to atmcmpxchg instructions are now emitted in the correct order Simulator · New Features o Addition of snull, stpick, ststream operations o The assembler now verifies that all scopes are properly closed o Stats added to track memory accesses that are not naturally aligned · Modified Features o Stream comparison operators changed to also return a control sequence: cmprelT cmpres.CRd.i1, valres.CRd.T, ctlseqres.CRd.i1, ctlseqa.CRLu.i1, vala.CRLu.T, ctlseqb.CRLu.i1, valb.CRLu.T · Removed Features - a number of deprecated operations have been removed from the simulator o Signed scalar loads (ldsN, ldsNd, ldsNx) have been removed. The unsigned scalar loads (ldN, ldNd, ldNx) are unchanged o All operations that used f34 or f66 datatypes have been removed · Pending Removal o The old “wide” memory references (ldNx8/stNx8) are deprecated and will be going away in the future. These should not be confused with the wide streaming memory operations (sldNx8/sstNx8), which are effectively superset replacements. o The seqc operation will be removed from the examples and the simulator as well. Code that uses them should be converted. o While they will not be removed, the atomics other than atmswp and atmcmpswp are not expected to be supported for V1 and should not be used for code. If you have a compelling example where these would make a significant performance difference to applications, please bring that to our attention. · Fixed Bugs o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected o Atmcmpxchg now returns the correct value · Known Issues o Atomic operations are modeled as a memory read o The fence operation is currently not invoking the memory model to wait for memory o The simple memory model is not waiting until memory quiesces before returning from a function. -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory

1 0

[Anl-intel-csa] Fwd: redaddfN
by Kumaran, Kalyan 28 Feb '18

28 Feb '18

Kent, Elliott, could one of you explain this and if this is the expected behavior in real hardware: o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. Thanks, Kumar Begin forwarded message: From: Vitali A Morozov <morozov(a)anl.gov<mailto:[email protected]>> Subject: Fwd: Re: redaddfN Date: February 27, 2018 at 11:16:13 AM CST To: "Kumaran, Kalyan" <kumaran(a)alcf.anl.gov<mailto:[email protected]>> Kumar, FYI. A short version of this thread: in a recent Feb 22 update, Intel has changed the behavior of a fp reduction unit. It is not pipelined anymore, accepting input once every 4 cycles. Most of us think this change will have significant impact on performance and should be inverted back. Thomas has implemented a temporal workaround to replace this unit with a pipelined graph at the expense of complexity and latency. Vitali -------- Forwarded Message -------- Subject: Re: redaddfN Date: Tue, 27 Feb 2018 10:45:30 -0600 From: Applencourt, Thomas <tapplencourt(a)anl.gov><mailto:[email protected]> To: Morozov, Vitali A. <morozov(a)alcf.anl.gov><mailto:[email protected]> Yes. You can find the implementation here: `/soft/restricted/intel_csa/csatools/csa-p/lib/libreduc.S` The implementation is straightforward. I use four readaddf32 in a round-robin fashion and do a manual reduction add the end. The only `trick`, is to duplicate the end signal to be able to stop all the 4 reduction unit. For that, I use the `replace` unit. To be totally honest, this is not a strict equivalent of the `readdf32` unit because I reorder the operation during the reduction, so numerical instability may occur. It's the price to pay to be able to pipeline. 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: Morozov, Vitali A. Sent: Tuesday, February 27, 2018 10:38 To: Applencourt, Thomas Subject: Re: redaddfN Hi Thomas, if I understand this thread correctly, you have implemented a pipelined reduction graph with the functionality, similar to readaddf32 unit. Is this correct? Vitali On 2/23/18 11:33 AM, Hal Finkel wrote: Hi, Thomas, Can you share the code? -Hal On 02/22/2018 04:32 PM, Applencourt, Thomas wrote: Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.)§ All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected This will have a big impact on application reeling on redaddf32N. So I did my own implementation [1] of this unit who allow pipelining like the old one. This new unit is a drop-in replacement of the old one. You can find the timing bellow: old : 4476 new_intel: 16678 new_anl : 4493 Regards, [1]: Hal, I dit it. I use one of their `strange` unit, the replace1 one. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 11:57 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version The discovery of new units continue. * completion (completion{8,16,32,64}availidx.Cd.u8, ldres.Cd.iN, stidx.Cu.u8, stdata.Cu.iN, size.Lu.u8) ``` Another mode of operation for scratchpads are as a “completion buffer”, where values can be inserted in a “sliding window”, and values can be retrieved in order. The completion operation supports this functionality. Note that this is a very unusual operator in terms of how it behaves. It can be viewed as effectively 3 independent operations bound together. I'm not yet understanding this unit. I have the feeling that it's a important one. - Snull (snull res.CRd.i1, ctlseq.CRLu.il<http://ctlseq.crlu.il/>) determine if current stream has zero length. Needed I think to implemented a more efficient `reduction add`. - Filter (filter{0-64} res.CRd.iN, filter.CRu.i1, val.CRLu.iN) This operation filters a set of values based on corresponding predicates. If filter is true, val is copied to res. If filter is false, val is consumed, but res is not written. syntactic sugar for switchN %ign, r, f, v - They also add a flag in the {pick/switch}any unit. This flag is used to change the mode of selection (priority, or round-robin). pickany{0-64} dst.CRd.iN, idx.CRd.i1, i0.CRLu.iN, i1.CRLu.iN, mode.Lu.i1=0 If the mode operand has the value 0 (default), this performs a prioritized pick, selecting i0 preferentially over i1 if both are available. If the mode operand has value 1, this will pick the last operand not selected, if available. ``` Implementing a Load balance strategy is now far more easy. 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: Applencourt, Thomas Sent: Thursday, February 22, 2018 09:19 To: Morozov, Vitali A.; Finkel, Hal J.; Knight, Christopher J. Cc: Bertoni, Colleen; Jin, Zheming Subject: Re: CSA SDK Release Notes for 2018/02/21 version Note that they add the documentation for 2 additional units: * prefetch (yeah!) * guard (sync a value with a signal) They also document new MEMLEVEL values (used by the memory ops): MEMLEVEL_T0 Use all levels of the cache hierarchy as appropriate MEMLEVEL_T1 Avoid any very local cache, if present, but use the main CSA cache. For CSA Version 1, there is no difference from using MEMLEVEL_T0 MEMLEVEL_T2 Bypass the first two levels of cache. For CSA Version 1, this implies avoiding the tile level cache. MEMLEVEL_NTA This asserts there is no benefit to caching (Non-Temporal Always), and references should bypass the (note: architectural cache hierarchy to the degree possible. NOTE: This is not just a hint, and things like memory ordering operands may not provide the same guarantees when this is used. 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: anl-intel-csa anl-intel-csa-bounces(a)lists.jlse.anl.gov<http://mailto:[email protected]/> on behalf of Glossop, Kent kent.glossop(a)intel.com<http://mailto:[email protected]/> Intel | Data Center Solutions, IoT, and PC Innovation<http://mailto:kent%[email protected]/> intel.com<http://intel.com/> Intel's innovation in cloud computing, data center, Internet of Things, and PC solutions is powering the smart and connected digital world we live in. Sent: Wednesday, February 21, 2018 22:54 To: Glossop, Kent Subject: [Anl-intel-csa] CSA SDK Release Notes for 2018/02/21 version CSA SDK Release Notes for 2018/02/21 version Intel Corporation Proprietary The default release has been changed to 2018-02-21-release. Relevant changes since the 2017-01-22-release: Documentation * The CSA Architecture document has finally been updated in the doc. area to reflect new operations and update operation signatures. It should be current with the simulator in this release. We expect to keep them synchronized in the future. The first entry in the edit history after the table of contents has a list of the additions/changes, and there are change bars in the margin. Spatial Advisor This includes a major update of Spatial Advisor which now incorporates per-application and per-loopnest offload strategy analyzer, L2 cache and HBM simple performance modeling and numerous other updates. New features: * Introduced analysis across overall program and all loopnests, not just innermost loops * Stabilized loopnest offload strategy selector * Per-program (Amdhal’s law) area and speed-up metrics introduced * L2 + HBM (external memory) performance model added; cache simulation enabled. * Performance estimate consideration of offload costs, such as configuration, invocation, and (optional) data transfer costs * New outputs, including a graphical representation of the call tree with offloaded, non-accelerable and offloadable regions highlighted * More precise area model for “Heavy INT” operations * Introduced “bound by” CSA workloads characterization to estimate fractions of PE-, memory-, cache- bound code. Compiler / Runtime * Inner loop Pipelining o Inner loop pipelining code generation now specifies minimum backedge LIC depths on the output assembly. This avoids deadlock even if the simulator creates shallow LICs by default. (Previously, ILPL assumed that all LICs would have a depth of 128.) o Inner loop pipelining code generation now waits for all of a gang of cohorts to be available before admitting them. This eliminates the possibility of deadlock in certain cases where buffering outside of the inner loop is insufficient. However, in these cases, performance will not be optimal due to the insufficient buffering. Future work on buffer insertion should address this. * SPMDization o Improved SPMDization intrinsic detection, resolving some crashes and incorrect code generation in cases involving fully unrolled loops. o Added a warning for loops where blocked SPMDization was requested but where the trip count cannot be determined by the compiler, which use to cause crashes. * Memory reference ordering o Fixed ordering of sequential outer loops containing parallel inner loops. o Fixed the handling of instructions that both read and write memory, which were previously mistakenly ordered as loads. o Re-implemented dropping of redundant dependencies from a particular memory instruction back to itself. o Removed old memory ordering passes and some of the old internal options for accessing them. At one point, we had recommended passing “-mllvm -csa-order-memops-type=wavefront” (one of the options that has been removed by this change) in order to avoid suboptimal handling of SPMDization by what was then the latest memory ordering pass. Since the current pass does not have that problem, this is no longer necessary and any uses of that option can be dropped. * Atomics o The compiler now emits atomics as “compare and swap” loops with atmcmpxchg rather than single “read modify write” operations such as “atmadd”. This better aligns with current hardware expectations. For experimentation purposes, the RMW-style operations can be restored by adding -Xclang -target-feature -Xclang ‘+rmwatomic’ to CFLAGS/CPPFLAGS/CSA_COMPILE_OPTIONS. Note that the simulator timing of atomic operations is still considered optimistic. o Operands to atmcmpxchg instructions are now emitted in the correct order Simulator * New Features o Addition of snull, stpick, ststream operations o The assembler now verifies that all scopes are properly closed o Stats added to track memory accesses that are not naturally aligned * Modified Features o Stream comparison operators changed to also return a control sequence: cmprelT cmpres.CRd.i1, valres.CRd.T, ctlseqres.CRd.i1, ctlseqa.CRLu.i1, vala.CRLu.T, ctlseqb.CRLu.i1, valb.CRLu.T * Removed Features - a number of deprecated operations have been removed from the simulator o Signed scalar loads (ldsN, ldsNd, ldsNx) have been removed. The unsigned scalar loads (ldN, ldNd, ldNx) are unchanged o All operations that used f34 or f66 datatypes have been removed * Pending Removal o The old “wide” memory references (ldNx8/stNx8) are deprecated and will be going away in the future. These should not be confused with the wide streaming memory operations (sldNx8/sstNx8), which are effectively superset replacements. o The seqc operation will be removed from the examples and the simulator as well. Code that uses them should be converted. o While they will not be removed, the atomics other than atmswp and atmcmpswp are not expected to be supported for V1 and should not be used for code. If you have a compelling example where these would make a significant performance difference to applications, please bring that to our attention. * Fixed Bugs o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) § All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. § redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers § The declaration of external output operand for memory operations has been corrected o Atmcmpxchg now returns the correct value * Known Issues o Atomic operations are modeled as a memory read o The fence operation is currently not invoking the memory model to wait for memory o The simple memory model is not waiting until memory quiesces before returning from a function. -- Hal Finkel Lead, Compiler Technology and Programming Languages Leadership Computing Facility Argonne National Laboratory

1 0

[Anl-intel-csa] Creation of the Bug, Issue and Request Tracker confluence page.
by Applencourt, Thomas 26 Feb '18

26 Feb '18

Hi, In order to keep track of the different ALCF feature requests / bug-reports, a confluence page was created: https://wiki.jlse.anl.gov/display/intelcsa/Bug%2C+Issue+and+Request+Tracker Please feel free to modify it, Regards, Thomas 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] floating-point latency and power
by Jin, Zheming 23 Feb '18

23 Feb '18

I made a simple experiment for the energy consumed for the floating-point vector add and divide program shown below. The experimental result shows that there are almost no differences in the "total energy measured ieu" for the floating-point add and divide operators. I assume a divide operator consumes more power. In addition, can you please explain that the latency is the same for all the floating-point operators ? Does a double-precision floating-point operator take longer latency than a single-precision floating-point operator ? Is it possible that the Intel architecture team can present a slide that shows the relationship between the power and latency for floating-point operators ? Thanks void example (int n, T * __restrict__ a, T * __restrict__ b, T * __restrict__ c) { int i; #pragma omp target map(to: a,b) map(from: c) { __builtin_csa_parallel_loop(); for (i=0; i<n; i++){ c[i] = a[i] / b[i]; // or add } } }

1 0

Re: [Anl-intel-csa] Conference call this Friday (02/23) at 1:00pm central
by Kazutomo Yoshii 23 Feb '18

23 Feb '18

If time allows, I'd like to discuss: * CSA driver model, kernel driver implementation, HW interfaces (control, debug, etc), APIs to userspace runtime I guess it may be too early to ask about some of these topics. It would be nice if Intel can share some overall idea, status, etc with us at this point. - kaz On 02/21/2018 08:42 PM, Kumaran, Kalyan wrote: > We will have our regularly scheduled conference call this Friday, Feb 23rd, at 1:00pm central. > > For Argonne folks I have reserved 4313. > > For folks dialing in you will find the Skype information on the confluence page (A21 page). > > Agenda: > > * Deep dive/discussion on the memory model supported in our simulator. See Thomas bug reports & questions (Elliott) > * Atomics supported on CSA (Elliott) > > -Kumar > > P.S.: Vitali and I will miss the call as we have to be on a call to discuss the SOW benchmark projections. Scott Parker will lead. > _______________________________________________ > anl-intel-csa mailing list > anl-intel-csa(a)lists.jlse.anl.gov > https://lists.jlse.anl.gov/mailman/listinfo/anl-intel-csa

1 0

[Anl-intel-csa] CSA SDK Release Notes for 2018/02/21 version
by Glossop, Kent 21 Feb '18

21 Feb '18

CSA SDK Release Notes for 2018/02/21 version Intel Corporation Proprietary The default release has been changed to 2018-02-21-release. Relevant changes since the 2017-01-22-release: Documentation - The CSA Architecture document has finally been updated in the doc. area to reflect new operations and update operation signatures. It should be current with the simulator in this release. We expect to keep them synchronized in the future. The first entry in the edit history after the table of contents has a list of the additions/changes, and there are change bars in the margin. Spatial Advisor This includes a major update of Spatial Advisor which now incorporates per-application and per-loopnest offload strategy analyzer, L2 cache and HBM simple performance modeling and numerous other updates. New features: - Introduced analysis across overall program and all loopnests, not just innermost loops - Stabilized loopnest offload strategy selector - Per-program (Amdhal's law) area and speed-up metrics introduced - L2 + HBM (external memory) performance model added; cache simulation enabled. - Performance estimate consideration of offload costs, such as configuration, invocation, and (optional) data transfer costs - New outputs, including a graphical representation of the call tree with offloaded, non-accelerable and offloadable regions highlighted - More precise area model for "Heavy INT" operations - Introduced "bound by" CSA workloads characterization to estimate fractions of PE-, memory-, cache- bound code. Compiler / Runtime - Inner loop Pipelining o Inner loop pipelining code generation now specifies minimum backedge LIC depths on the output assembly. This avoids deadlock even if the simulator creates shallow LICs by default. (Previously, ILPL assumed that all LICs would have a depth of 128.) o Inner loop pipelining code generation now waits for all of a gang of cohorts to be available before admitting them. This eliminates the possibility of deadlock in certain cases where buffering outside of the inner loop is insufficient. However, in these cases, performance will not be optimal due to the insufficient buffering. Future work on buffer insertion should address this. - SPMDization o Improved SPMDization intrinsic detection, resolving some crashes and incorrect code generation in cases involving fully unrolled loops. o Added a warning for loops where blocked SPMDization was requested but where the trip count cannot be determined by the compiler, which use to cause crashes. - Memory reference ordering o Fixed ordering of sequential outer loops containing parallel inner loops. o Fixed the handling of instructions that both read and write memory, which were previously mistakenly ordered as loads. o Re-implemented dropping of redundant dependencies from a particular memory instruction back to itself. o Removed old memory ordering passes and some of the old internal options for accessing them. At one point, we had recommended passing "-mllvm -csa-order-memops-type=wavefront" (one of the options that has been removed by this change) in order to avoid suboptimal handling of SPMDization by what was then the latest memory ordering pass. Since the current pass does not have that problem, this is no longer necessary and any uses of that option can be dropped. - Atomics o The compiler now emits atomics as "compare and swap" loops with atmcmpxchg rather than single "read modify write" operations such as "atmadd". This better aligns with current hardware expectations. For experimentation purposes, the RMW-style operations can be restored by adding -Xclang -target-feature -Xclang '+rmwatomic' to CFLAGS/CPPFLAGS/CSA_COMPILE_OPTIONS. Note that the simulator timing of atomic operations is still considered optimistic. o Operands to atmcmpxchg instructions are now emitted in the correct order Simulator - New Features o Addition of snull, stpick, ststream operations o The assembler now verifies that all scopes are properly closed o Stats added to track memory accesses that are not naturally aligned - Modified Features o Stream comparison operators changed to also return a control sequence: cmprelT cmpres.CRd.i1, valres.CRd.T, ctlseqres.CRd.i1, ctlseqa.CRLu.i1, vala.CRLu.T, ctlseqb.CRLu.i1, valb.CRLu.T - Removed Features - a number of deprecated operations have been removed from the simulator o Signed scalar loads (ldsN, ldsNd, ldsNx) have been removed. The unsigned scalar loads (ldN, ldNd, ldNx) are unchanged o All operations that used f34 or f66 datatypes have been removed - Pending Removal o The old "wide" memory references (ldNx8/stNx8) are deprecated and will be going away in the future. These should not be confused with the wide streaming memory operations (sldNx8/sstNx8), which are effectively superset replacements. o The seqc operation will be removed from the examples and the simulator as well. Code that uses them should be converted. o While they will not be removed, the atomics other than atmswp and atmcmpswp are not expected to be supported for V1 and should not be used for code. If you have a compelling example where these would make a significant performance difference to applications, please bring that to our attention. - Fixed Bugs o Floating point reducer latency has changed, which may require recoding of reductions to avoid a significant performance hit. (Often it will be desirable to pipeline multiple partial sums through a normal operation that can accept a new value per cycle, and reduce the partial sums at the end, rather than using a reduction operator.) ? All floating point reducers now have a latency of 4 cycles, and only accept a value once every 4 cycles. ? redaddfN, redsubfN, sredaddfN, sredsubfN latency was erroneously 1 cycle. It is now 4 cycles like the rest of the floating point reducers ? The declaration of external output operand for memory operations has been corrected o Atmcmpxchg now returns the correct value - Known Issues o Atomic operations are modeled as a memory read o The fence operation is currently not invoking the memory model to wait for memory o The simple memory model is not waiting until memory quiesces before returning from a function.

1 0