https://bugs.freedesktop.org/show_bug.cgi?id=60879
--- Comment #37 from Hristo Venev mustrumr97@gmail.com --- I updated llvm, clang and mesa. Division and modulo still don't work. Another thing I noticed is that ifs which depend on memory loads cause llvm crash:
__kernel void add(__global const uint *a, __global const uint *b, __global uint *c){ ulong id=get_global_id(0); // OK if(id>10) return; // OK if(b[id]==0) return; // crash c[id]=a[id]/b[id]; // GPU hang }
a[id] is id+1 b[id] is 2*id+2
Stack dump: 0. Running pass 'Function Pass Manager' on module 'radeon'. 1. Running pass 'AMDGPU DAG->DAG Pattern Instruction Selection' on function '@add' Segmentation fault
#0 0x00007ffff461c8a7 in ?? () from /usr/lib64/llvm/libLLVM-3.4svn.so #1 0x00007ffff3e36208 in llvm::SelectionDAGISel::DoInstructionSelection() () from /usr/lib64/llvm/libLLVM-3.4svn.so #2 0x00007ffff3e3c620 in llvm::SelectionDAGISel::CodeGenAndEmitDAG() () from /usr/lib64/llvm/libLLVM-3.4svn.so #3 0x00007ffff3e3e0f2 in llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) () from /usr/lib64/llvm/libLLVM-3.4svn.so #4 0x00007ffff3e3f421 in llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) () from /usr/lib64/llvm/libLLVM-3.4svn.so #5 0x00007ffff3acaeb2 in llvm::FPPassManager::runOnFunction(llvm::Function&) () from /usr/lib64/llvm/libLLVM-3.4svn.so #6 0x00007ffff3acaf4b in llvm::FPPassManager::runOnModule(llvm::Module&) () from /usr/lib64/llvm/libLLVM-3.4svn.so #7 0x00007ffff3acb195 in llvm::MPPassManager::runOnModule(llvm::Module&) () from /usr/lib64/llvm/libLLVM-3.4svn.so #8 0x00007ffff3acd1dc in llvm::PassManagerImpl::run(llvm::Module&) () from /usr/lib64/llvm/libLLVM-3.4svn.so #9 0x00007ffff417c009 in ?? () from /usr/lib64/llvm/libLLVM-3.4svn.so #10 0x00007ffff417c382 in LLVMTargetMachineEmitToMemoryBuffer () from /usr/lib64/llvm/libLLVM-3.4svn.so #11 0x00007ffff2ae6ab1 in radeon_llvm_compile () from /usr/lib64/gallium-pipe/pipe_radeonsi.so #12 0x00007ffff2adc65d in si_compile_llvm () from /usr/lib64/gallium-pipe/pipe_radeonsi.so #13 0x00007ffff2adef79 in ?? () from /usr/lib64/gallium-pipe/pipe_radeonsi.so #14 0x00007ffff6d882a7 in _cl_kernel::exec_context::bind(_cl_command_queue*) () from /usr/lib64/libOpenCL.so.1 #15 0x00007ffff6d88e46 in _cl_kernel::launch(_cl_command_queue&, std::vector<unsigned long, std::allocator<unsigned long> > const&, std::vector<unsigned long, std::allocator<unsigned long> > const&, std::vector<unsigned long, std::allocator<unsigned long> > const&) () from /usr/lib64/libOpenCL.so.1 #16 0x00007ffff6d847dc in _cl_event::trigger() () from /usr/lib64/libOpenCL.so.1 #17 0x00007ffff6d84e54 in clover::hard_event::hard_event(_cl_command_queue&, unsigned int, std::vector<_cl_event*, std::allocator<_cl_event*> >, std::function<void (_cl_event&)>) () from /usr/lib64/libOpenCL.so.1 #18 0x00007ffff6d9fad5 in clEnqueueNDRangeKernel () from /usr/lib64/libOpenCL.so.1