Displaying 5 results from an estimated 5 matches for "siinstruct".
Did you mean:
instruct
2013 Dec 31
2
[LLVMdev] [PATCH] R600 - Fix zero extend of i1
Hi,
When trying to compile a trivial opencl kernel such as:
__kernel void if_eq(__global int * out, int arg0, int arg1){
out[0] = arg0==arg1?0:1;
}
Clang generates IR like:
%1 = icmp eq i32 %arg0, %arg1
%. = zext i1 %1 to i32
This eventually crashes ISel on R600. Attached patch adds a selector so
it will compile.
Regards,
Jon Pry
jonpry at gmail.com
-------------- next
2016 Mar 28
0
RFC: atomic operations on SI+
...rget/AMDGPU/AMDGPUISelLowering.h | 1 +
> lib/Target/AMDGPU/CIInstructions.td | 3 +-
> lib/Target/AMDGPU/SIISelLowering.cpp | 41 +++++++++++++++
> lib/Target/AMDGPU/SIISelLowering.h | 1 +
> lib/Target/AMDGPU/SIInstrInfo.td | 9 ++++
> lib/Target/AMDGPU/SIInstructions.td | 2 +-
> test/CodeGen/AMDGPU/global_atomics.ll | 89 ++++++++++++++++++++++++++++++++
> 8 files changed, 145 insertions(+), 2 deletions(-)
>
> diff --git a/lib/Target/AMDGPU/AMDGPUISelLowering.cpp b/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
> index 0820898..6c5b1ed 1...
2016 Mar 25
2
RFC: atomic operations on SI+
Hi Tom, Matt,
I'm working on a project that needs few coherent atomic operations (HSA
mode: load, store, compare-and-swap) for std::atomic_uint in HCC.
the attached patch implements atomic compare and swap for SI+
(untested). I tried to stay within what was available, but there are
few issues that I was unsure how to address:
1.) it currently uses v2i32 for both input and output. This
2020 Feb 19
2
i1 true ^= -1 in DAG matcher?
...1 true, i1 true, i1 true, i1 true, i1 true, i1
true, i1 true, i1 true>
[use of %7]
results in vector of '-1' in the DAG. This also seems the reason why
LLVM's vnot PatFrag doesn't match in this case. I've also found f
rom third_party/llvm/llvm-project/llvm/lib/Target/AMDGPU/SIInstructions.td:
// FIXME: The generated DAG matcher seems to have strange behavior
// with a 1-bit literal to match, so use a -1 for checking a true
// 1-bit value.
Which seems to reflect the same observation. Is this a bug or a feature?
Thanks in advance for any explanation
-------------- next part ---...
2012 Jul 16
3
[LLVMdev] RFC: LLVM incubation, or requirements for committing new backends
...llvm/trunk/lib/Target/AMDGPU/SIISelLowering.h
> llvm/trunk/lib/Target/AMDGPU/SIInstrFormats.td
> llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp
> llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.h
> llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
> llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
> llvm/trunk/lib/Target/AMDGPU/SIIntrinsics.td
> llvm/trunk/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
> llvm/trunk/lib/Target/AMDGPU/SIMachineFunctionInfo.h
> llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.cpp
> llvm/trunk/lib/Target/AMDGPU/SIRegisterInfo.h...