January 2023 - Ecp-aurora-sdk-discuss

Re: [Ecp-aurora-sdk-discuss] [EXTERNAL] who disables fast math (informal poll)?
by Bryce Allen 24 Jan '23

24 Jan '23

My main interest here is what applications are using for production science runs, not trying to rock the boat on what should/shouldn't be the default. I haven't found any papers written that try to capture this (I guess it is challenging with so many applications and different machines and people running them). Some anecdotes is better than nothing, thanks for the replies. For the SyclCPLX library, I think we've landed on fast-math enabled => most aggressive implementation, not enabled => minimize error. This came up with the division operator, which is one of the most expensive when doing infinite and NaN checks and can also be implemented with some clever scaling to minimize error. But this all comes at a cost, and the existing implementation for std::complex in oneAPI does not do scaling, so we are going to follow that to avoid regressing performance. -Bryce On 1/24/23 17:54, Kent, Paul via Ecp-aurora-sdk-discuss wrote: > I think it will be difficult to get any consensus here. There are plenty > of historical discussions around different compiler defaults, and > different communities have different reasons to argue one way or another. > > In our case we ensure all our unit and integration tests will pass under > aggressive optimization since that reflects the production compilation > of our app, and we know that aggressive optimization helps get the best > performance out of the hardware. Test failures are usually seen as > indicating that the app should not be used for production science. > > -- Paul Kent > > *From: *Ecp-aurora-sdk-discuss > <ecp-aurora-sdk-discuss-bounces(a)lists.jlse.anl.gov> on behalf of Arndt, > Daniel via Ecp-aurora-sdk-discuss > <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> > *Date: *Tuesday, January 24, 2023 at 4:30 PM > *To: *Bryce Allen <ballen(a)anl.gov>, > ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov > <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> > *Subject: *Re: [Ecp-aurora-sdk-discuss] [EXTERNAL] who disables fast > math (informal poll)? > > I always use `-fp-model=precise` for unit testing. There are a bunch of > unit tests in projects I care about (Kokkos-based) that would fail > otherwise. I think it’s a very unfortunate choice to choose > (floating-point) optimization flags that aggressive as default. > > Best, > > Daniel > > *From: *Ecp-aurora-sdk-discuss > <ecp-aurora-sdk-discuss-bounces(a)lists.jlse.anl.gov> on behalf of Bryce > Allen via Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> > *Date: *Tuesday, January 24, 2023 at 1:11 PM > *To: *ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov > <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> > *Subject: *[EXTERNAL] [Ecp-aurora-sdk-discuss] who disables fast math > (informal poll)? > > The default for dpcpp / icpx is to enable fast math. How many people > here disable this because they need the extra accuracy? What is the > default for other compilers when using OpenMP target offload? Anyone > know of survey studies that look at this across different types of > applications? > > This has come up in the context of standardizing complex numbers in > sycl, by adding sycl::complex: > https://gcc02.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.co… <https://gcc02.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.co…> > > Thanks, > Bryce > >

1 0

Re: [Ecp-aurora-sdk-discuss] [EXTERNAL] who disables fast math (informal poll)?
by Kent, Paul 24 Jan '23

24 Jan '23

I think it will be difficult to get any consensus here. There are plenty of historical discussions around different compiler defaults, and different communities have different reasons to argue one way or another. In our case we ensure all our unit and integration tests will pass under aggressive optimization since that reflects the production compilation of our app, and we know that aggressive optimization helps get the best performance out of the hardware. Test failures are usually seen as indicating that the app should not be used for production science. -- Paul Kent From: Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss-bounces(a)lists.jlse.anl.gov> on behalf of Arndt, Daniel via Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Date: Tuesday, January 24, 2023 at 4:30 PM To: Bryce Allen <ballen(a)anl.gov>, ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: Re: [Ecp-aurora-sdk-discuss] [EXTERNAL] who disables fast math (informal poll)? I always use `-fp-model=precise` for unit testing. There are a bunch of unit tests in projects I care about (Kokkos-based) that would fail otherwise. I think it’s a very unfortunate choice to choose (floating-point) optimization flags that aggressive as default. Best, Daniel From: Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss-bounces(a)lists.jlse.anl.gov> on behalf of Bryce Allen via Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Date: Tuesday, January 24, 2023 at 1:11 PM To: ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: [EXTERNAL] [Ecp-aurora-sdk-discuss] who disables fast math (informal poll)? The default for dpcpp / icpx is to enable fast math. How many people here disable this because they need the extra accuracy? What is the default for other compilers when using OpenMP target offload? Anyone know of survey studies that look at this across different types of applications? This has come up in the context of standardizing complex numbers in sycl, by adding sycl::complex: https://gcc02.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.co…<https://gcc02.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.co…> Thanks, Bryce

1 0

Re: [Ecp-aurora-sdk-discuss] [EXTERNAL] who disables fast math (informal poll)?
by Arndt, Daniel 24 Jan '23

24 Jan '23

I always use `-fp-model=precise` for unit testing. There are a bunch of unit tests in projects I care about (Kokkos-based) that would fail otherwise. I think it’s a very unfortunate choice to choose (floating-point) optimization flags that aggressive as default. Best, Daniel From: Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss-bounces(a)lists.jlse.anl.gov> on behalf of Bryce Allen via Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Date: Tuesday, January 24, 2023 at 1:11 PM To: ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: [EXTERNAL] [Ecp-aurora-sdk-discuss] who disables fast math (informal poll)? The default for dpcpp / icpx is to enable fast math. How many people here disable this because they need the extra accuracy? What is the default for other compilers when using OpenMP target offload? Anyone know of survey studies that look at this across different types of applications? This has come up in the context of standardizing complex numbers in sycl, by adding sycl::complex: https://gcc02.safelinks.protection.outlook.com/?url=https%3A%2F%2Fgithub.co… Thanks, Bryce

1 0

who disables fast math (informal poll)?
by Bryce Allen 24 Jan '23

24 Jan '23

The default for dpcpp / icpx is to enable fast math. How many people here disable this because they need the extra accuracy? What is the default for other compilers when using OpenMP target offload? Anyone know of survey studies that look at this across different types of applications? This has come up in the context of standardizing complex numbers in sycl, by adding sycl::complex: https://github.com/argonne-lcf/SyclCPLX/pull/29 Thanks, Bryce

1 0

florentia_debug nodes moving to SUSE SP4 Image
by Muralidharan, Servesh 23 Jan '23

23 Jan '23

Hi florentia_debug node users, We will be upgrading the node image on the nodes in florentia_debug queue to a new image based on the OpenSUSE SP4 image. This includes migrating to a much newer Linux Kernel and major update to the GPU KMD stack. Once we test this image out for a week or two, the rest of the florentia nodes in the main queue will be upgraded to the same image as well. Please let us know if you run into any new issues on these nodes. On behalf of early hw team. -- Servesh

1 0

Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions
by Luo, Ye 12 Jan '23

12 Jan '23

Hi Patrick, I read your reply of this email thread again. I think the transfer you referred is not relevant to QMCPACK since I only do intra team reduction. The reduction buffer in global memory is relevant. I assume the hanging I observed was just caused by bugs when handling the pool of buffers. Do you have updates if CMPLRLLVM-42440 has been fixed or not? Ye ________________________________ From: Luo, Ye <yeluo(a)anl.gov> Sent: Wednesday, December 7, 2022 7:19 AM To: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com>; ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions The performance of reduction is not my concern. Reduction is only the last part of a kernel. Even if it takes 2x time, the overall kernel time increases by 10%. Let me reiterate that the key problem is the transfer. Every action added to interaction to the GPU adds contention. I can only accept using the buffer implementation once it has no additional transfer. So looking forward to a fix removing the artificial transfer. Ye ________________________________ From: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com> Sent: Tuesday, December 6, 2022 12:36 PM To: Luo, Ye <yeluo(a)anl.gov>; ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Cc: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions Thomas, can you add this as a topic to compiler meeting? >Not sure how you worded CMPLRLLVM-42440. My request is not only not hang but also no added transfer for reduction pointers. It is to fix the hang. The current plan and implementation of buffer approach has not and will not have any explicit data transfers. Based on unit testing I expect to see about a 3x reduction in latency compared to standard approach without any buffers. Results of course depend on what you actually do, we can take a look at performance of your case once we have fixed the hang. From: Luo, Ye <yeluo(a)anl.gov> Date: Tuesday, December 6, 2022 at 12:14 To: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com>, ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions It is not about what I think. SLM is in fact better now as I prove in my cases and buffer is terrible for now. Please fix the issue and show me the buffer is competitive in my test case, then I may agree that buffer scheme is good enough. We cannot say a broken implementation is better. Not sure how you worded CMPLRLLVM-42440. My request is not only not hang but also no added transfer for reduction pointers. Ye ________________________________ From: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com> Sent: Tuesday, December 6, 2022 11:23 AM To: Luo, Ye <yeluo(a)anl.gov>; ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Cc: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions The plan is to just have buffer approach. While you might think the SLM approach is better for stability, given your example, it is actually the other way around. The buffer approach has more unit tests passing. The difficulty is that there is no knowledge on how much SLM you have available during generation of spv kernels. Only IGC knows about it. I don’t think it is worth it to get an automatic switch between buffer and SLM as there is no performance benefit. For now you can use that compiler flag as a workaround while we work on fixing your hang under CMPLRLLVM-42440. From: Luo, Ye <yeluo(a)anl.gov> Date: Tuesday, December 6, 2022 at 10:25 To: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com>, ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions I finished testing QMCPACK "ctest -R deterministic". -mllvm -vpo-paropt-atomic-free-reduction-slm=true removes all my headaches. 100% tests passed, 0 tests failed out of 719 Total Test time (real) = 625.22 sec This is at the same level of NVIDIA and AMD timing. The follow-up question is how to expose this option or make it more automated. As long as SLM is not exhausted, I would prefer using SLM rather than global memory and stay away from runtime challenges. Ye ________________________________ From: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com> Sent: Monday, December 5, 2022 10:17 AM To: Luo, Ye <yeluo(a)anl.gov>; ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Cc: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions >He can just build the QMCPACK and run the determinsitic tests "ctest -R deterministic". Can you put the reproducer in the ANL bug tracker? How many threads do you have running in parallel doing a reduction? I know that we have a similar issue when we do a reduction with 16 threads host threads per stack in parallel but I haven’t looked into it yet as it is not a commend scheme. If you want I can file my reproducer and tag it as QMCPACK internally. > atomic on the global memory? Hopefully it is atomic to the L1/SLM memory only. In our case no atomics are used on the global buffer. Just a barrier. >1024 * 16 * 8 byte is only 128KB. L1/SLM has 512KB. So what is the default split of SLM and L1? SLM is limited to 128 not 512. >Considering all the vector engines within a XE core can only do 256 Ops per cycle, I probably prefer 256 to 512 workgroup size rather than oversubscribing a lot. In my performance unit tests, the performance is better with 1024 team size. > Could you share the compiler option? I am not sure what features you have in your compiler version but you may try: -mllvm -vpo-paropt-atomic-free-reduction-slm=true and -mllvm -vpo-paropt-atomic-free-reduction-slm=false and see if the hang goes away. From: Luo, Ye <yeluo(a)anl.gov> Date: Monday, December 5, 2022 at 09:38 To: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com>, ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions Unlocking a single lock on host is fast. It should already be used in your compiler version to get a global buffer pool since it is not allocated. I should have said the additional thread contention instead of locking/unlocking. > Not doing cross-team reduction. I need a reduction with threads within each team work very well. In this case, atomic should never been used. A team consists of multiple EU hardware-threads which requires the use of synchronization and/or use of atomic. atomic on the global memory? Hopefully it is atomic to the L1/SLM memory only. > I feel 100x more than necessary memory is consumed. I think this comes from minimal L0 allocation size? Note sure. > Has this been implemented and tested? If yes, do you have a rough ETA? If not yet, is there a backup plan if this proposed sceme doesn't work? I'm wondering if we can have an implementation using SLM and not worry about n-th buffer, most of the reduction I need is just a few scalars or fixed size array no more than 16 doubles, Even for your reduction with 16 doubles it would not work with SLM. You typically want to run the reduction with largest team size (1024). So the minimal amount of SLM you would need is: 1024 * 16 * 8 byte which is larger than SLM size. 1024 * 16 * 8 byte is only 128KB. L1/SLM has 512KB. So what is the default split of SLM and L1? Considering all the vector engines within a XE core can only do 256 Ops per cycle, I probably prefer 256 to 512 workgroup size rather than oversubscribing a lot. Also my actual kernel only need reduction of 4 doubles. The buffers are not expected to add overhead so I would not worry about performance. Especially in your case as you are just doing a team local reduction which uses only the team global buffer. There is a compiler switch to turn it off. We have a lot of reduction unit tests. Could you share the compiler option? Ye From: Luo, Ye <yeluo(a)anl.gov> Date: Sunday, December 4, 2022 at 23:15 To: Steinbrecher, Patrick <patrick.steinbrecher(a)intel.com>, ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: Re: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions Intel compiler uses a global buffer for OpenMP reductions. Then each team offsets into that global buffer to a unique chunk to perform the local team reduction. We have studied this extensively and concluded this is the fastest way for doing the reduction. This local chunk will just be cached in L1. There is no need to use SLM as it makes no performance difference. This gives more SLM space to the users kernel. Also, this enables larger reductions to work without running out of SLM. That global buffer will use a memory pool that can be re-used. No allocation will happen in reductions (except maybe the first call). Probably I can live with this. Not sure if you take into account the cost of locks when allocating/deallocating scratch from a pool in multi threaded codes. To enable atomic free reductions we also added a 2nd global buffer of size 1 that is zero initialized that is used to figure which is the last team running. This also stems from a zero initialized pool, no allocations should be visible during multiple calls of reductions. Not doing cross-team reduction. I need a reduction with threads within each team work very well. In this case, atomic should never been used. In one of the next releases, we will add a 3rd buffer to get rid of all explicit data transfers needed for reduction variable map(tofrom). Has this been implemented and tested? If yes, do you have a rough ETA? If not yet, is there a backup plan if this proposed sceme doesn't work? I'm wondering if we can have an implementation using SLM and not worry about n-th buffer, most of the reduction I need is just a few scalars or fixed size array no more than 16 doubles, Right now reduction on 4 scalars. Libomptarget --> Obtained target argument (Begin: 0xff00fffffffd0a40, Offset: 0) from host pointer 0x0000150719c06740 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002400000 for private array 0x0000000000000000 - pushing target argument 0xff00000002400000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002000000 for private array 0x0000000000000000 - pushing target argument 0xff00000002000000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002800000 for private array 0x0000000000000000 - pushing target argument 0xff00000002800000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002c00000 for private array 0x0000000000000000 - pushing target argument 0xff00000002c00000 Target LEVEL0 RTL --> Team sizes = {64, 1, 1} Target LEVEL0 RTL --> Number of teams = {2, 1, 1} I feel 100x more than necessary memory is consumed. All of the above will enable reductions with a latency equivalent to an empty kernel call. with the added lock/unlock, I doubt "equivalent" is true 😉. Ye From: Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss-bounces(a)lists.jlse.anl.gov> on behalf of Luo, Ye via Ecp-aurora-sdk-discuss <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Date: Friday, December 2, 2022 at 23:10 To: ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov <ecp-aurora-sdk-discuss(a)lists.jlse.anl.gov> Subject: [Ecp-aurora-sdk-discuss] A few suspeciouis memory operation around kernel launch in OpenMP target regions Hi compiler engineers, I'm looking at the execution of this kernel https://github.com/QMCPACK/qmcpack/blob/3c3b92cede781a02eb33c68e10ff118e5ea… 1. with LIBOMPTARGET_DEBUG, I saw Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002400000 for private array 0x0000000000000000 - pushing target argument 0xff00000002400000 My kernel only do reduction within each team. With upstream clang, to my best knowledge, there is no need of any global memory and using shared/local memory is sufficient. Any working in progress to get rid of the need of such scratch pad in the global memory? 1. When I intercept the code with gdb and I frequently saw PrivateArgumentManagerTy been invoked with allocating and deleting memory (see backtrace at the end). but my code has only first private pointers and scalars and doesn't need data transfer. Is it because the reduction scratch pointers are transferred via the PrivateArgumentManagerTy mechanism? Trying to minimize libomptarget interaction with the plugin. Ye LIBOMPTARGET_DEBUG info Libomptarget --> Device 0 is ready to use. Libomptarget --> Entry 0: Base=0x000014a9cb139000, Begin=0x000014a9cb139000, Size=512, Type=0x21, Name=mw_dist[:dist_stride * nw] Libomptarget --> Entry 1: Base=0x000014a9dba0f680, Begin=0x000014a9dba0f680, Size=32, Type=0x25, Name=transfer_buffer_ptr[:transfer_buffer.size()] Libomptarget --> Entry 2: Base=0x000014a9caff6600, Begin=0x000014a9caff6600, Size=384, Type=0x22, Name=mw_cur_allu[:n_padded * 3 * nw] Libomptarget --> Entry 3: Base=0x000014a9dba0f200, Begin=0x000014a9dba0f200, Size=16, Type=0x21, Name=grp_ids[:n_src] Libomptarget --> Entry 4: Base=0x000014a9dba0f640, Begin=0x000014a9dba0f640, Size=40, Type=0x26, Name=mw_vgl[:(DIM + 2) * nw] Libomptarget --> Entry 5: Base=0x0000000000000000, Begin=0x0000000000000000, Size=4194304, Type=0x4a0, Name=unknown Libomptarget --> Entry 6: Base=0x0000000000000000, Begin=0x0000000000000000, Size=4194304, Type=0x4a0, Name=unknown Libomptarget --> Entry 7: Base=0x0000000000000000, Begin=0x0000000000000000, Size=4194304, Type=0x4a0, Name=unknown Libomptarget --> Entry 8: Base=0x0000000000000000, Begin=0x0000000000000000, Size=4194304, Type=0x4a0, Name=unknown Libomptarget --> Entry 9: Base=0x0000000000000000, Begin=0x0000000000000000, Size=4194304, Type=0x4a0, Name=unknown Libomptarget --> Entry 10: Base=0x0000000000000040, Begin=0x0000000000000040, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 11: Base=0x0000000000000000, Begin=0x0000000000000000, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 12: Base=0x0000000000000001, Begin=0x0000000000000001, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 13: Base=0x0000000000000002, Begin=0x0000000000000002, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 14: Base=0x0000000000000001, Begin=0x0000000000000001, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 15: Base=0x0000000000000002, Begin=0x0000000000000002, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 16: Base=0x0000000000000004, Begin=0x0000000000000004, Size=0, Type=0x120, Name=unknown Libomptarget --> Entry 17: Base=0x0000000000000010, Begin=0x0000000000000010, Size=0, Type=0x120, Name=unknown Libomptarget --> loop trip count is 0. Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9cb139000, Size=512)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9cb139000, TgtPtrBegin=0xff00fffffff91000, Size=512, DynRefCount=2 (incremented), HoldRefCount=0, Name=mw_dist[:dist_stride * nw] Libomptarget --> There are 512 bytes allocated at target address 0xff00fffffff91000 - is not new Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9dba0f680, Size=32)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9dba0f680, TgtPtrBegin=0xff00fffffffd08c0, Size=32, DynRefCount=2 (incremented), HoldRefCount=0, Name=transfer_buffer_ptr[:transfer_buffer.size()] Libomptarget --> Moving 32 bytes (hst:0x000014a9dba0f680) -> (tgt:0xff00fffffffd08c0) Target LEVEL0 RTL --> Copied 32 bytes (hst:0x000014a9dba0f680) -> (tgt:0xff00fffffffd08c0) Libomptarget --> There are 32 bytes allocated at target address 0xff00fffffffd08c0 - is not new Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9caff6600, Size=384)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9caff6600, TgtPtrBegin=0xff00fffffffc0e00, Size=384, DynRefCount=2 (incremented), HoldRefCount=0, Name=mw_cur_allu[:n_padded * 3 * nw] Libomptarget --> There are 384 bytes allocated at target address 0xff00fffffffc0e00 - is not new Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9dba0f200, Size=16)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9dba0f200, TgtPtrBegin=0xff00fffffffd0380, Size=16, DynRefCount=2 (incremented), HoldRefCount=0, Name=grp_ids[:n_src] Libomptarget --> There are 16 bytes allocated at target address 0xff00fffffffd0380 - is not new Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9dba0f640, Size=40)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9dba0f640, TgtPtrBegin=0xff00fffffffd0880, Size=40, DynRefCount=2 (incremented), HoldRefCount=0, Name=mw_vgl[:(DIM + 2) * nw] Libomptarget --> There are 40 bytes allocated at target address 0xff00fffffffd0880 - is not new Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9cb139000, Size=512)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9cb139000, TgtPtrBegin=0xff00fffffff91000, Size=512, DynRefCount=2 (update suppressed), HoldRefCount=0 Libomptarget --> Obtained target argument (Begin: 0xff00fffffff91000, Offset: 0) from host pointer 0x000014a9cb139000 Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9dba0f680, Size=32)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9dba0f680, TgtPtrBegin=0xff00fffffffd08c0, Size=32, DynRefCount=2 (update suppressed), HoldRefCount=0 Libomptarget --> Obtained target argument (Begin: 0xff00fffffffd08c0, Offset: 0) from host pointer 0x000014a9dba0f680 Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9caff6600, Size=384)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9caff6600, TgtPtrBegin=0xff00fffffffc0e00, Size=384, DynRefCount=2 (update suppressed), HoldRefCount=0 Libomptarget --> Obtained target argument (Begin: 0xff00fffffffc0e00, Offset: 0) from host pointer 0x000014a9caff6600 Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9dba0f200, Size=16)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9dba0f200, TgtPtrBegin=0xff00fffffffd0380, Size=16, DynRefCount=2 (update suppressed), HoldRefCount=0 Libomptarget --> Obtained target argument (Begin: 0xff00fffffffd0380, Offset: 0) from host pointer 0x000014a9dba0f200 Libomptarget --> Looking up mapping(HstPtrBegin=0x000014a9dba0f640, Size=40)... Libomptarget --> Mapping exists with HstPtrBegin=0x000014a9dba0f640, TgtPtrBegin=0xff00fffffffd0880, Size=40, DynRefCount=2 (update suppressed), HoldRefCount=0 Libomptarget --> Obtained target argument (Begin: 0xff00fffffffd0880, Offset: 0) from host pointer 0x000014a9dba0f640 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002400000 for private array 0x0000000000000000 - pushing target argument 0xff00000002400000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002000000 for private array 0x0000000000000000 - pushing target argument 0xff00000002000000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002800000 for private array 0x0000000000000000 - pushing target argument 0xff00000002800000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000002c00000 for private array 0x0000000000000000 - pushing target argument 0xff00000002c00000 Target LEVEL0 RTL --> Allocated 4194304 bytes from reduction scratch pool Libomptarget --> Allocated 4194304 bytes of target memory at 0xff00000003000000 for private array 0x0000000000000000 - pushing target argument 0xff00000003000000 Libomptarget --> Forwarding first-private value 0x0000000000000040 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000000 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000001 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000002 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000001 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000002 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000004 to the target construct Libomptarget --> Forwarding first-private value 0x0000000000000010 to the target construct Libomptarget --> Launching target execution __omp_offloading_bf36f4fa_c000da9f__ZN11qmcplusplus14BsplineFunctorIfE14mw_evaluateVGLEiiPKPKS1_iPKiiPfiPKfS8_RNS_6VectorIcNS_12OMPallocatorIcNS_17SYCLHostAllocatorIcLm64EEEEEEE_l59 with pointer 0x0000000009c89690 (index=0). Target LEVEL0 RTL --> Executing a kernel 0x0000000009c89690... Target LEVEL0 RTL --> Team sizes = {1024, 1, 1} Target LEVEL0 RTL --> Number of teams = {448, 1, 1} Target LEVEL0 RTL --> Kernel Pointer argument 0 (value: 0xff00fffffff91000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 1 (value: 0xff00fffffffd08c0) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 2 (value: 0xff00fffffffc0e00) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 3 (value: 0xff00fffffffd0380) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 4 (value: 0xff00fffffffd0880) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 5 (value: 0xff00000002400000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 6 (value: 0xff00000002000000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 7 (value: 0xff00000002800000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 8 (value: 0xff00000002c00000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Pointer argument 9 (value: 0xff00000003000000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 10 (value: 0x0000000000000040) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 11 (value: 0x0000000000000000) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 12 (value: 0x0000000000000001) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 13 (value: 0x0000000000000002) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 14 (value: 0x0000000000000001) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 15 (value: 0x0000000000000002) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 16 (value: 0x0000000000000004) was set successfully for device 0. Target LEVEL0 RTL --> Kernel Scalar argument 17 (value: 0x0000000000000010) was set successfully for device 0. Target LEVEL0 RTL --> Submitted kernel 0x0000000004870b90 to device 0 Target LEVEL0 RTL --> Executed kernel entry 0x0000000009c89690 on device 0 the backtrace shows MemAllocatorTy::dealloc tries to free device memory. #0 0x0000146eac4d0463 in __lll_lock_wait () from /lib64/libpthread.so.0 #1 0x0000146eac4c8251 in pthread_mutex_lock () from /lib64/libpthread.so.0 #2 0x0000146ea82c38f1 in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #3 0x0000146ea823199d in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #4 0x0000146ea8231f6d in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #5 0x0000146ea82c3eab in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #6 0x0000146ea8343b25 in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #7 0x0000146ea8343570 in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #8 0x0000146ea834e428 in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #9 0x0000146ea834fcbb in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #10 0x0000146ea834ff7f in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #11 0x0000146ea80455c1 in ?? () from /soft/libraries/intel-gpu-umd/20221031.1-pvc-prq-66/driver/lib64/libze_intel_gpu.so.1 #12 0x0000146ea8f935f9 in MemAllocatorTy::dealloc (this=this@entry=0x35fc408, Ptr=Ptr@entry=0xff00fffffd000000) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/plugins/level0/src/rtl.cpp:2790 #13 0x0000146ea8f8d612 in RTLDeviceInfoTy::dataDelete (this=<optimized out>, DeviceId=<optimized out>, Ptr=0xff00fffffd000000) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/plugins/level0/src/rtl.cpp:4860 #14 0x0000146eac9603ce in DeviceTy::deleteData (this=0x36a3790, TgtPtrBegin=0xff00fffffd000000, Kind=3) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/src/device.cpp:623 #15 0x0000146eac973987 in (anonymous namespace)::PrivateArgumentManagerTy::free (this=0x7ffe6e3fadf0) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/src/omptarget.cpp:1552 #16 (anonymous namespace)::processDataAfter (Loc=0x15beb70, DeviceId=0, ArgNum=18, ArgBases=0x7ffe6e3fb228, Args=0x7ffe6e3fb198, PrivateArgumentManager=..., HostPtr=<optimized out>, ArgSizes=<optimized out>, ArgTypes=<optimized out>, ArgNames=<optimized out>, ArgMappers=<optimized out>, AsyncInfo=...) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/src/omptarget.cpp:1839 #17 target (Loc=Loc@entry=0x15beb70, Device=..., HostPtr=HostPtr@entry=0xc5b45f <__omp_offloading_bf36f4fa_c000da9f__ZN11qmcplusplus14BsplineFunctorIfE14mw_evaluateVGLEiiPKPKS1_iPKiiPfiPKfS8_RNS_6VectorIcNS_12OMPallocatorIcNS_17SYCLHostAllocatorIcLm64EEEEEEE_l59.region_id>, ArgNum=18, ArgBases=0x7ffe6e3fb228, Args=0x7ffe6e3fb198, ArgSizes=<optimized out>, ArgTypes=<optimized out>, ArgNames=<optimized out>, ArgMappers=<optimized out>, TeamNum=<optimized out>, ThreadLimit=<optimized out>, Tripcount=<optimized out>, IsTeamConstruct=<optimized out>, AsyncInfo=...) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/src/omptarget.cpp:1983 #18 0x0000146eac965e47 in __tgt_target_kernel (Loc=<optimized out>, DeviceId=<optimized out>, NumTeams=-512, ThreadLimit=<optimized out>, HostPtr=<optimized out>, Args=0x7ffe6e3fb0f0) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/src/interface.cpp:393 #19 0x0000146eac9819fa in __tgt_target_teams_mapper (Loc=0x35fda40, DeviceId=128, HostPtr=0x282, ArgNum=<optimized out>, ArgsBase=<optimized out>, Args=<optimized out>, ArgSizes=0x7ffe6e3fb2b8, ArgTypes=0xc5b460, ArgNames=0xc5b4f0, ArgMappers=0x0, NumTeams=0, ThreadLimit=0) at /netbatch/donb843633_00/runDir/dir/workspace/NIT/xmain/LX/xmainefi2linux_release/ws/icsws/llvm/openmp/libomptarget/src/LegacyAPI.cpp:124 #20 0x0000000000a20654 in qmcplusplus::BsplineFunctor<float>::mw_evaluateVGL (iat=<optimized out>, num_groups=<optimized out>, functors=<optimized out>, n_src=16, grp_ids=0x146ea0938240, nw=<optimized out>, mw_vgl=<optimized out>, n_padded=<optimized out>, mw_dist=<optimized out>, mw_cur_allu=<optimized out>, transfer_buffer=...) at /home/yeluo/opt/qmcpack/src/QMCWaveFunctions/Jastrow/BsplineFunctor.cpp:59 ... =================== Ye Luo, Ph.D. Computational Science Division & Leadership Computing Facility Argonne National Laboratory

1 0