User Details
- User Since
- Oct 29 2014, 9:58 AM (437 w, 5 d)
Today
I wonder if you need to update getAliasResult in AMDGPUAliasAnalysis for the new address space.
Fri, Mar 17
From the description:
llvm.is.fpclass(x, fcPosInf) -> fcmp oeq x, +inf
llvm.is.fpclass(x, fcNegInf) -> fcmp oeq x, -inf
llvm.is.fpclass(x, ~fcPosInf) -> fcmp one x, +inf
llvm.is.fpclass(x, ~fcNegInf) -> fcmp one x, -inf
The last two should be "une".
Patch seems fine to me, but I wonder if we need to get more eyes on your SetVector changes. I don't think there is a code owner for ADT. @dblaikie?
Thu, Mar 16
Wed, Mar 15
Thanks!
Tue, Mar 14
LGTM.
Fold (bitcast (logicop (bitcast x), (c))) -> (and x, (bitcast c)) iff the current logicop type is illegal
Looks good with nits.
The commit message should make it clear that you're also changing it to use UniformityAnalysis instead of (Legacy?)DivergenceAnalysis.
LGTM but I'd like @sameerds to have the final say.
Looks good overall.
Mon, Mar 13
Fri, Mar 10
I'm not competent to review the implementation but it seems to work well for me. This will let us replace lots of instances of foreach _ = BoolToList<cond>.ret in ... with if cond then ... in AMDGPU backend tablegen files.
Thu, Mar 9
but I think this also ends up fixing the original reported issue (https://github.com/llvm/llvm-project/issues/49830) as well
Wed, Mar 8
LGTM, thanks.
Tue, Mar 7
Hi, this patch makes tablegen crash instead of error on cases like this:
if !cast<A>("").x then def x;
(Based on the test case for https://github.com/llvm/llvm-project/issues/49830)
LGTM.
Looks OK to me, thanks.
Just my 2p: it feels a bit premature to commit patches for this. It feels more like something you could prototype on a branch somewhere and come back when you have more experience with how it all works out in practice.
Mon, Mar 6
Update
Thanks for the fix! Unfortunately I am not competent to review it.
LGTM, thanks!
Sun, Mar 5
Fri, Mar 3
Works for me, thanks! I have a larger test case that goes from about 50 s to less than 1 ms with this patch.
Thu, Mar 2
Thanks!
Is there a reason this is implemented in C++ instead of instruction selection patterns?
Hi, this patch is provoking a build failure in an LLVM_ENABLE_EXPENSIVE_CHECKS build on my Ubuntu 22.04.2 machine using Clang 14 as the host compiler:
[189/189] Building CXX object unittests/ADT/CMakeFiles/ADTTests.dir/STLExtrasTest.cpp.o FAILED: unittests/ADT/CMakeFiles/ADTTests.dir/STLExtrasTest.cpp.o /usr/lib/ccache/clang++ -DEXPENSIVE_CHECKS -DGTEST_HAS_RTTI=0 -D_DEBUG -D_GLIBCXX_ASSERTIONS -D_GLIBCXX_DEBUG -D_GNU_SOURCE -D_LIBCPP_ENABLE_ASSERTIONS -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -I/home/jayfoad2/llvm-expensive/unittests/ADT -I/home/jayfoad2/git/llvm-project/llvm/unittests/ADT -I/home/jayfoad2/llvm-expensive/include -I/home/jayfoad2/git/llvm-project/llvm/include -I/home/jayfoad2/git/llvm-project/third-party/unittest/googletest/include -I/home/jayfoad2/git/llvm-project/third-party/unittest/googlemock/include -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wc++98-compat-extra-semi -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wsuggest-override -Wstring-conversion -Wmisleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -g -Wno-variadic-macros -Wno-gnu-zero-variadic-macro-arguments -fno-exceptions -fno-rtti -Wno-suggest-override -std=c++17 -MD -MT unittests/ADT/CMakeFiles/ADTTests.dir/STLExtrasTest.cpp.o -MF unittests/ADT/CMakeFiles/ADTTests.dir/STLExtrasTest.cpp.o.d -o unittests/ADT/CMakeFiles/ADTTests.dir/STLExtrasTest.cpp.o -c /home/jayfoad2/git/llvm-project/llvm/unittests/ADT/STLExtrasTest.cpp In file included from /home/jayfoad2/git/llvm-project/llvm/unittests/ADT/STLExtrasTest.cpp:9: In file included from /home/jayfoad2/git/llvm-project/llvm/include/llvm/ADT/STLExtras.h:20: In file included from /home/jayfoad2/git/llvm-project/llvm/include/llvm/ADT/Hashing.h:51: In file included from /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/algorithm:60: In file included from /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/bits/stl_algobase.h:69: In file included from /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/debug.h:133: /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/functions.h:110:44: error: no matching function for call to '__addressof' return __foreign_iterator_aux4(__it, std::__addressof(*__other)); ^~~~~~~~~~~~~~~~ /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/functions.h:164:14: note: in instantiation of function template specialization '__gnu_debug::__foreign_iterator_aux3<__gnu_cxx::__normal_iterator<const int *, std::__cxx1998::vector<int>>, std::vector<int>, std::random_access_iterator_tag, llvm::detail::SafeIntIterator<int, false>>' requested here return __foreign_iterator_aux3(__it, __other, __other_end, __tag()); ^ /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/functions.h:186:5: note: in instantiation of function template specialization '__gnu_debug::__foreign_iterator_aux2<__gnu_cxx::__normal_iterator<const int *, std::__cxx1998::vector<int>>, std::vector<int>, std::random_access_iterator_tag, llvm::detail::SafeIntIterator<int, false>>' requested here || __foreign_iterator_aux2(__it, std::__miter_base(__other), ^ /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/functions.h:198:14: note: in instantiation of function template specialization '__gnu_debug::__foreign_iterator_aux<__gnu_cxx::__normal_iterator<const int *, std::__cxx1998::vector<int>>, std::vector<int>, std::random_access_iterator_tag, llvm::detail::SafeIntIterator<int, false>>' requested here return __foreign_iterator_aux(__it, __other, __other_end, _Integral()); ^ /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/vector:626:4: note: in instantiation of function template specialization '__gnu_debug::__foreign_iterator<__gnu_cxx::__normal_iterator<const int *, std::__cxx1998::vector<int>>, std::vector<int>, std::random_access_iterator_tag, llvm::detail::SafeIntIterator<int, false>>' requested here __glibcxx_check_insert_range(__position, __first, __last, __dist); ^ /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/debug/macros.h:183:36: note: expanded from macro '__glibcxx_check_insert_range' _GLIBCXX_DEBUG_VERIFY(__gnu_debug::__foreign_iterator(_Position,_First,_Last),\ ^ /home/jayfoad2/git/llvm-project/llvm/include/llvm/ADT/STLExtras.h:2017:5: note: in instantiation of function template specialization 'std::vector<int>::insert<llvm::detail::SafeIntIterator<int, false>, void>' requested here C.insert(C.end(), adl_begin(R), adl_end(R)); ^ /home/jayfoad2/git/llvm-project/llvm/unittests/ADT/STLExtrasTest.cpp:367:3: note: in instantiation of function template specialization 'llvm::append_range<std::vector<int>, llvm::iota_range<int>>' requested here append_range(V, llvm::seq(6, 8)); ^ /usr/bin/../lib/gcc/x86_64-linux-gnu/12/../../../../include/c++/12/bits/move.h:49:5: note: candidate function [with _Tp = int] not viable: expects an lvalue for 1st argument __addressof(_Tp& __r) _GLIBCXX_NOEXCEPT ^ 1 error generated. ninja: build stopped: subcommand failed.
Do you have any idea what is wrong here? Thanks!
Looks good, thanks!
Wed, Mar 1
LGTM.
LGTM.
LGTM.
LGTM, thanks.
Tue, Feb 28
Seems obvious
I don't quite understand the bug here - can you explain how "s_load_dword s1, s[2:3], s0 0x1" would have been mishandled before your patch?
The 0x1 bit is currently indistinguishable from offset:0x1 and therefore accepted as a valid operand.
I'm fine with the patch but I don't particularly like the fact that DA has an isDivergent(Instruction *) in the first place, because it's really unclear what it means - for non-void instructions I guess it's the same as isDivergent(Value *), and for conditional branches it means something specific, and for other instructions...? I'd much prefer for have a more specific isDivergentBranch(BranchInst *).
I don't quite understand the bug here - can you explain how "s_load_dword s1, s[2:3], s0 0x1" would have been mishandled before your patch?
This is because it introduces a need to have a build_vector before copying the PHI value, and that build_vector may have many undef elements. This can cause very high register pressure and abnormal stack usage in some cases.