similar to: [LLVMdev] [NVPTX] Backend failure in LegalizeDAG due to unimplemented expand in target lowering

Hi again, Kind people on #llvm helped me to utilize bugpoint to reduce the previously submitted test case. For record, it code be done with the following command: $ bugpoint -llc-safe test.ll The resulting IR is attached, and it is crashing in the same way. Is it a valid code? dmikushin at hp2:~/forge/kernelgen/branches/tests_lnt/behavior/sincos> llc test.ll.1 This action is not supported

Dear colleagues, I'm looking if we can replace nvopencc with LLVM NVPTX in our project. It turns NVPTX won't work with the code nvopencc can handle (please see the log below). So are atomic intrinsics not supported or am I doing call in a wrong way? Thanks, - Dima. SOURCE ======== dmikushin at hp2:~> cat kernelgen_monitor.ll ; ModuleID =

> -----Original Message----- > From: Dmitry N. Mikushin [mailto:maemarcus at gmail.com] > Sent: Wednesday, May 16, 2012 5:44 AM > To: LLVM-Dev > Cc: Justin Holewinski > Subject: NVPTX: __iAtomicCAS support ? > > Dear colleagues, > > I'm looking if we can replace nvopencc with LLVM NVPTX in our project. > It turns NVPTX won't work with the code nvopencc

Hi all, considering following .ll file ; ModuleID = 'vect3x.ll' target triple = "armv7-none-linux-gnueabi" define arm_aapcscc void @test_hi_char8(i8* %.T0351, <8 x i8>* nocapture %srcA, <4 x i8>* nocapture %dst) noinline { L.entry: %0 = tail call arm_aapcscc i32 (...)* @get_global_id(i8* %.T0351, i32 0) %1 = bitcast <8 x i8>* %srcA to <4 x i8>*

Sounds like a bug in vector promote. If I restore this flag and use -promote-elements=0 everything works for me. Please fill a PR in LLVM bugzilla and assign to Nadav. On Mon, Jun 25, 2012 at 5:04 PM, Sebastien DELDON-GNB <sebastien.deldon at st.com> wrote: > Hi all, > > > considering following .ll file > > ; ModuleID = 'vect3x.ll' > target triple =

Dear all, Looks like the NVPTX backend cannot handle array-of-arrays contant (please see the reporocase below). Is it supposed to work? Any ideas how to get it working? Important for our target applications. Thanks, - Dima. $ cat test.ll ; ModuleID = '__kernelgen_main_module' target datalayout =

I believe this patch has broken a PPC app that I am tracking. Here is a reduced test case. Reproduce with llc -mattr=+Altivec -mcpu=g5. The backtrace looks like this: #0 0x9333ae42 in __kill () #1 0x9333ae34 in kill$UNIX2003 () #2 0x933ad23a in raise () #3 0x933b9679 in abort () #4 0x933ae3db in __assert_rtn () #5 0x0008bd8f in llvm::MVT::getVectorElementType (this=0xbfffdda4) at

I think our test case demonstrates that requiring the array item being initialized to be constant is incorrect. NVPTX does not crash anymore and produces correct result with the following change: --- NVPTXAsmPrinter.cpp 2012-09-03 15:14:00.000000000 +0200 +++ NVPTXAsmPrinter.cpp 2012-09-04 15:47:17.859398193 +0200 @@ -1890,17 +1890,15 @@ case Type::ArrayTyID: case Type::VectorTyID: case

Things are still broken. Unfortunately llvm test suite does not contain enough vector code to fully test this. Can you revert the patch first? Evan On Feb 24, 2009, at 7:14 PM, Scott Michel wrote: > Evan: > > I did not encounter this back trace before I committed the newest > BuildVectorSDNode patch, which removed all class instance members > and passes results back via

NVCC successfully handles the same IR, if we try to process the same .cu file with clang+nvptx and nvcc: CLANG/NVPTX: ============= $ cat dayofweek.cu __attribute__((device)) char yweek[7][4] = { "MON", "TUE", "WED", "THU", "FRI", "SAT", "SUN" }; $ clang -cc1 -emit-llvm -fcuda-is-device dayofweek.cu -o dayofweek.ll $ cat

Hi, I'm trying to use llvm/clang to cross compile to sparcv9. The following works with a -march=sparc, but yields errors for sparcv9. Are there some other flags that need to be specified? Thanks, Tarun > clang -m64 -emit-llvm test.c -c -o test.bc > llc -march=sparcv9 test.bc -o hello.s ExpandIntegerResult #0: 0x8a6c478: i64 = GlobalAddress<[4 x i8]* @.str> 0 [ORD=1] [ID=0] Do

Evan: I work on reverting it, although, when I tried yesterday, it wasn't particularly clean (lots of rejected patches, presumably due to intervening commits.) Are you still getting the backtrace or is this just a case of incorrectly generated code? -scooter On Wed, Feb 25, 2009 at 10:09 AM, Evan Cheng <echeng at apple.com> wrote: > Things are still broken. Unfortunately llvm

On 09/04/2012 09:57 AM, Dmitry N. Mikushin wrote: > I think our test case demonstrates that requiring the array item being > initialized to be constant is incorrect. NVPTX does not crash anymore > and produces correct result with the following change: > > --- NVPTXAsmPrinter.cpp 2012-09-03 15:14:00.000000000 +0200 > +++ NVPTXAsmPrinter.cpp 2012-09-04 15:47:17.859398193 +0200 >

Evan: I did not encounter this back trace before I committed the newest BuildVectorSDNode patch, which removed all class instance members and passes results back via reference parameters. -scooter On Tue, Feb 24, 2009 at 11:39 AM, Evan Cheng <evan.cheng at apple.com> wrote: > I believe this patch has broken a PPC app that I am tracking. Here is a > reduced test case. Reproduce

Scott, In case you missed it, I reimplemented your BuildVectorSDNode::isConstantSplat method following the suggestions from Chris. The revised version passes "make check" for llvm. Assuming that it also passes Evan's tests, I think it should also do what you need for CellSPU. On Feb 25, 2009, at 12:16 PM, Scott Michel wrote: > Evan: > > I work on reverting it,

Duncan: I'm still stymied how this whole thread ended up about shuffle vector nodes, when the original problem was my build vector patch. I'm still working on backing the build vector patch out (it isn't clean with all of the intervening commits and I have pressing management tasks which command my attention.) -scooter On Tue, Feb 24, 2009 at 12:28 AM, Duncan Sands <baldrick at

> 3. Introduce a new ShuffleVectorSDNode that only has two SDValue > operands (the two input vectors), but that also contains an array of > ints in the node (not as operands). ... > The important part of #3 is that we really want an array of ints > (using -1 for undef) for the shuffle mask, not "operands". This > eliminates the nastiness we have now were we

Dear Zhang, Compiler ends up invoking cc1 (the backend) anyways. So if you would like to invoke it by hand, the only thing to know is the right combination of options. Try to use the compiler verbose option "-v". It will show you how exactly clang invokes the backend: > clang -v -c showdebug.c clang version 3.2 (trunk 156703) Target: x86_64-unknown-linux-gnu Thread model: posix

Per subject, this patch adding an additional pass to handle vector operations; the idea is that this allows removing the code from LegalizeDAG that handles illegal types, which should be a significant simplification. There are still some issues with this patch, but does the approach look sane? -Eli -------------- next part -------------- Index: lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp

Hi, After building out project in release mode, caught an assertion, which we have not seen before: hello_f: /tmp/rpmbuild_debug/BUILD/llvm/build/include/llvm/ADT/DenseMap.h:126: void llvm::DenseMap<KeyT, ValueT, KeyInfoT>::clear() [with KeyT = llvm::MachineBasicBlock*, ValueT = <unnamed>::BlockChain*, KeyInfoT = llvm::DenseMapInfo<llvm::MachineBasicBlock*>]: Assertion