search for: siinstructions

Displaying 5 results from an estimated 5 matches for "siinstructions".

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 10064...
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 >...