- Nov 11, 2023
-
-
Nikita Popov authored
This reverts commit c1d5865a. Introduces a new use of ConstantExpr::getAShr().
-
- Nov 10, 2023
-
-
Nikita Popov authored
Use the constant folding API instead. As we're working on ImmConstants, these folds are guaranteed to succeed.
-
Guray Ozen authored
This PR adds a test that performs GEMM 128x128x128 (F32 += F16 * F16). It uses `sm_90a` features in NVGPU dialect. Simplified algorithm is as follows: **Prologue** ``` mgroup = mbarriers.init x 2 tma.load ... shmem_buffer_lhs<0 x 128 x 64> tma.load ... shmem_buffer_rhs<0 x 64 x 64> tma.load ... shmem_buffer_rhs<0 x 64 x 64> mbarrier.expect_tx 32768 tma.load ... shmem_buffer_lhs<1 x 128 x 64> tma.load ... shmem_buffer_rhs<1 x 64 x 64> tma.load ... shmem_buffer_rhs<1 x 64 x 64> mbarrier.expect_tx 32768 ``` **Mainloop** ``` matrixD = for(i = 0;...2) { mbarrier.try_wait [i] lhs = shmem_buffer_lhs<pipe x 128 x 64> rhs = shmem_buffer_rhs<pipe x 64 x 128> yield nvgpu.warpgroup.mma (lhs, rhs) // Expanded : nvgpu.warpgroup.mma [128][128]+=[128][64]*[64][128] // wgmma.m64n128k16(A[0:64][0:16] * B[0:16][0:128]) // wgmma.m64n128k16(A[0:64][16:32] * B[16:32][0:128]) // wgmma.m64n128k16(A[0:64][32:48] * B[32:48][0:128]) // wgmma.m64n128k16(A[0:64][48:64] * B[48:64][0:128]) // wgmma.m64n128k16(A[64:128][0:16] * B[0:16][0:128]) // wgmma.m64n128k16(A[64:128][16:32] * B[16:32][0:128]) // wgmma.m64n128k16(A[64:128][32:48] * B[32:48][0:128]) // wgmma.m64n128k16(A[64:128][48:64] * B[48:64][0:128]) ``` **Epilogue** ``` //reg->shmem warpgroup.mma.store matrixD, shmem //shmem->glbmem parallel-for(i=0;...128) parallel-for(j=0;...128) store shmem, globalmem ``` -
Guray Ozen authored
PR #69913 added a GEMM test (128x128x128 F32 += F16 * F16) with if-statement. This PR adds the same test using predicates in PTX. Predicate support is enabled using _BasicPtxBuilderInterface_ `(nvgpu.opcode ..., predicate = %pred)`. The predicate condition is computed in `Step 2. [GPU] Elect fastest thread in CTA` inspired by cutlass. It is as follows: ``` lane_predicate = nvvm.elect.sync warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0) warp_idx_in_warp_group = warp_idx % 4 predicate = (lane_predicate & warp_idx_in_warp_group) ``` Depends on #70027 #69934 #69935 #69584
-
Anna Thomas authored
-
Guray Ozen authored
#70923 improved verifier. The verifier caught that the tensor map type in the tma descriptor in this test isn't correct. The program was working correctly anway since the offset is calculated correctly. This work fixes the test.
-
Guillaume Chatelet authored
The update is automatically generated from `config/config.json`.
-
Joseph Huber authored
Summary: The AMDGPU backend uses the linker-provided INIT_ARRAY and FINI_ARRAY sections to call all the global constructors in a single kernel. Previously this mistakenly used the same iteration logic for both arrays. The destructors stored in FINI_ARRAY are stored in the same order as the ones in the INIT_ARRAY section so we need to traverse it in reverse order.
-
Joseph Huber authored
Summary: This pass emits the new "nvptx$device$init" and "nvptx$device$fini" kernels that are callable by the device. This intends to mimic the method of lowering for AMDGPU where we emit `amdgcn.device.init` and `amdgcn.device.fini` respectively. These kernels simply iterate a symbol called `__init_array_start/stop` and `__fini_array_start/stop`. Normally, the linker provides these symbols automatically. In the AMDGPU case we only need call the kernel and we call the ctors / dtors. However, for NVPTX we require the user initializes these variables to the associated globals that we already emit as a part of this pass. The motivation behind this change is to move away from OpenMP's handling of ctors / dtors. I would much prefer that the backend / runtime handles this. That allows us to handle ctors / dtors in a language agnostic way, This approach requires that the runtime initializes the associated globals. They are marked `weak` so we can emit this per-TU. The kernel itself is `weak_odr` as it is copied exactly. One downside is that any module containing these kernels elicitis the "stack size cannot be statically determined warning" every time from `nvlink` which is annoying but inconsequential for functionality. It would be nice if there were a way to silence this warning however.
-
Nikita Popov authored
-
Nikita Popov authored
These will no longer be created by default during constant folding.
-
Nikita Popov authored
-
Nikita Popov authored
Otherwise we risk infinite loops when shift constant expressions are no longer supported.
-
Quinn Dawkins authored
This is the last step needed for basic support for distributing masked vector code. The lane id gets delinearized based on the distributed mask shape and then compared against the original mask sizes to compute the bounds for the distributed mask. Note that the distribution of masks is implicit on the shape specified by the warp op. As a result, it is the responsibility of the consumer of the mask to ensure the distributed mask will match its own distribution semantics.
-
Jan Svoboda authored
Following up on #69975, this patch skips writing `DIAG_PRAGMA_MAPPINGS` as well. Deserialization of this PCM record is still showing up in profiles, since it needs to be VBR-decoded for every transitively loaded PCM file. The scanner doesn't make any guarantees about diagnostic accuracy (and it even disables all warnings), so skipping this record should be safe.
-
Nikita Popov authored
Use the constant folding API instead.
-
Nikita Popov authored
Use the constant folding API instead. As we're working on ImmConstant, it is guaranteed to succeed.
-
diggerlin authored
copy implement of https://github.com/llvm/llvm-project/pull/71359 to AIX.cpp and add test scenario --------- Co-authored-by:
zhijian <zhijian@ca.ibm.com>
-
Sam Tebbs authored
Adds the following SME2 builtins: svmin_single_(s8|s16|s32|s64)_x(2|4) svmin_single_(u8|u16|u32|u64)_x(2|4) svmin_single_(f16|f32|f64)_x(2|4) svmin_(s8|s16|s32|s64)_x(2|4) svmin_(u8|u16|u32|u64)_x(2|4) svmin_(f16|f32|f64)_x(2|4) svmax_single_(s8|s16|s32|s64)_x(2|4) svmax_single_(u8|u16|u32|u64)_x(2|4) svmax_single_(f16|f32|f64)_x(2|4) svmax_(s8|s16|s32|s64)_x(2|4) svmax_(u8|u16|u32|u64)_x(2|4) svmax_(f16|f32|f64)_x(2|4) See [https://github.com/ARM-software/acle/pull/217](https://github.com/ARM-software/acle/pull/217) Patch by: Kerry McLaughlin <kerry.mclaughlin@arm.com> -
Quinn Dawkins authored
Following up on https://github.com/llvm/llvm-project/pull/71676#discussion_r1387561079
-
Nikita Popov authored
The complex set of type checks in this code reduces down to "always return nullptr". Drop the code to use the default implementation instead, which will just compute the KnownBits for the bitcast.
-
Jake Egan authored
We now have 64-bit XCOFF object file support, so these tests can be enabled again. However, some tests still fail due to unsupported debug sections, so I cleaned up their comments.
-
Nikita Popov authored
Work on APInt instead.
-
Ivan Kosarev authored
-
Ivan Kosarev authored
CPol and CPol_GLC1 operand classes have identical predicates, which means AsmParser cannot differentiate between the RTN and non-RTN variants of the instructions. When it currently selects the wrong instruction, a hack in cvtSMEMAtomic() corrects the op-code. Using the new predicated-value operands makes this hack and the whole conversion function not needed. Other uses of CPol_GLC1 operands are to be addressed separately. Resolves about half of the remaining ~1000 pairs of ambiguous instructions. Part of <https://github.com/llvm/llvm-project/issues/69256>.
-
Florian Mayer authored
-
Nikita Popov authored
The destination element type here cannot be a pointer type, as this would require an inttoptr rather than bitcast.
-
Quinn Dawkins authored
Currently when there is a mix of transfer read ops and transfer write ops that need to be distributed, because the pattern for write distribution is rooted on the transfer write, it is hard to guarantee that the write gets distributed after the read when the two aren't directly connected by SSA. This is likely still relatively unsafe when there are undistributable ops, but structurally these patterns are a bit difficult to work with. For now pattern benefits give fairly good guarantees for happy paths.
-
David Spickett authored
This reverts commit 343eb4b4. This test should work now that the build script doesn't add -m(32|64) on Arm/AArch64 builders.
-
Hans Wennborg authored
This causes asserts to fire: llvm/lib/Analysis/ValueTracking.cpp:4262: std::tuple<Value *, FPClassTest, FPClassTest> llvm::fcmpImpliesClass(CmpInst::Predicate, const Function &, Value *, const APFloat *, bool): Assertion `(RHSClass == fcPosNormal || RHSClass == fcNegNormal || RHSClass == fcPosSubnormal || RHSClass == fcNegSubnormal) && "should have been recognized as an exact class test"' failed. See comments on the PR. > Previously we could recognize exact class tests performed by > an fcmp with special values (0s, infs and smallest normal). > Expand this to recognize the implied classes by a compare with a general > constant. e.g. fcmp ogt x, 1 implies positive and non-0. > > The API should be better merged with fcmpToClassTest but that > made the diff way bigger, will try to do that in a future > patch. This reverts commit dc3faf0e.
-
David Green authored
Similar to #70635, this expands the handling of integer to fp conversions. The code is very similar to the float->integer conversions with types handled oppositely. There are some extra unhandled cases which require more handling for ASR operations.
-
Yingwei Zheng authored
This patch lowers `sdiv x, +/-2**k` to `add + select + shift` when the short forward branch optimization is enabled. The latter inst seq performs faster than the seq generated by target-independent DAGCombiner. This algorithm is described in ***Hacker's Delight***. This patch also removes duplicate logic in the X86 and AArch64 backend. But we cannot do this for the PowerPC backend since it generates a special instruction `addze`.
-
David Spickett authored
This reverts commit fd5206cc. Fixing the test case would require some awkard %if use that I'm not sure would even work, or splitting it into 2 copies that are almost identical. Instead, always add -m for clang, which allows it for all targets, but not for gcc which does not.
-
Eric authored
There are three changes present in this PR. 1. Use github packages for libcxx-builder rather than dockerhub. The ldionne/libcxx-builder image will now be hosted at ghcr.io/libcxx/libcxx-builder. This has the benefit of allowing members of the github org to push new versions. In the future I hope to add github actions to rebuild the image as needed. 2. Add docker-compose file The compose file allows to to specify the package repository, so that users can simply write 'docker compose build' and 'docker compose push'. It also gives us a centralized place to manage version arguments which change frequently. 3. Use non-shell CMD form. This may help the google libcxx builders disconnect more gracefully as the shell form of CMD may eat the shutdown signal. I'm hoping this corrects inaccurate agent counts from the buildkite API, since when the VM's terminate, they do so without signaling it to buildkite, whi...
-
NimishMishra authored
This patch skips default privatization for crashing cases like namelists, reduction instrinsics, and structure constructor. Fixes: https://github.com/llvm/llvm-project/issues/67332, https://github.com/llvm/llvm-project/issues/66454, and https://github.com/llvm/llvm-project/issues/65569 Co-Authored-By:
kiranchandramohan <kiran.chandramohan@arm.com>
-
David Spickett authored
This reverts commit 7c3603e1. Turns out when you ask for "clang" it also uses the GCC Builder class, so I need to update some tests.
-
David Spickett authored
-
Kadir Cetinkaya authored
-
David Spickett authored
This option is definitely needed for x86_64, and is valid for PowerPC and s390x too. I'm using "in" because on Armv8 Linux the uname is actually "armv8l" not just "arm".
-
Haojian Wu authored
The test added causes some buildbot failures: - https://lab.llvm.org/buildbot/#/builders/17/builds/45077 https://lab.llvm.org/buildbot/#/builders/96/builds/48277 Remove the test for now to make the builtbots happy, and will re-add it after investigation.
-