[llvm] r201222 - R600: Always implement both versions of isTruncateFree and add a sanity check.
Tom Stellard
tom at stellard.net
Thu Feb 13 08:06:01 PST 2014
On Thu, Feb 13, 2014 at 02:31:52AM +0000, Daenzer, Michel wrote:
> On Mit, 2014-02-12 at 10:17 +0000, Benjamin Kramer wrote:
> > Author: d0k
> > Date: Wed Feb 12 04:17:54 2014
> > New Revision: 201222
> >
> > URL: http://llvm.org/viewvc/llvm-project?rev=201222&view=rev
> > Log:
> > R600: Always implement both versions of isTruncateFree and add a sanity check.
>
> This change broke the piglit OpenCL test
> Program/Execute/gegl-rgb-gamma-u8-to-ragabaf with the radeonsi driver.
> I'm attaching the test output including a dump of the LLVM IR and a gdb
> backtrace.
>
This patch only indirectly caused the regression. The real culprit is
actually r200947. See this thread:
http://lists.cs.uiuc.edu/pipermail/llvm-commits/Week-of-Mon-20140210/204708.html
-Tom
>
> --
> Earthling Michel Dänzer | http://www.amd.com
> Libre software enthusiast | Mesa and X developer
> Program received signal SIGABRT, Aborted.
> 0x00007ffff6d201d5 in __GI_raise (sig=sig at entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
> 56 ../nptl/sysdeps/unix/sysv/linux/raise.c: No such file or directory.
> (gdb) bt
> #0 0x00007ffff6d201d5 in __GI_raise (sig=sig at entry=6) at ../nptl/sysdeps/unix/sysv/linux/raise.c:56
> #1 0x00007ffff6d23388 in __GI_abort () at abort.c:90
> #2 0x00007ffff233207c in llvm::llvm_unreachable_internal (msg=msg at entry=0x7ffff27d2e1f "Undefined function",
> file=file at entry=0x7ffff27d3198 "/home/daenzer/src/llvm-git/llvm/lib/Target/R600/AMDGPUISelLowering.h", line=line at entry=96)
> at /home/daenzer/src/llvm-git/llvm/lib/Support/ErrorHandling.cpp:99
> #3 0x00007ffff203aa55 in llvm::AMDGPUTargetLowering::LowerCall (this=<optimized out>, CLI=..., InVals=...) at /home/daenzer/src/llvm-git/llvm/lib/Target/R600/AMDGPUISelLowering.h:96
> #4 0x00007ffff22ac773 in llvm::TargetLowering::LowerCallTo (this=0x9d4440, CLI=...) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp:7229
> #5 0x00007ffff21ef12f in (anonymous namespace)::SelectionDAGLegalize::ExpandLibCall (LC=llvm::RTLIB::MUL_I64, Node=0xf35db0, isSigned=<optimized out>, this=<optimized out>,
> this=<optimized out>) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:1928
> #6 0x00007ffff21ef2e6 in (anonymous namespace)::SelectionDAGLegalize::ExpandIntLibCall (this=<optimized out>, Node=<optimized out>, isSigned=<optimized out>, Call_I8=<optimized out>,
> Call_I16=<optimized out>, Call_I32=<optimized out>, Call_I64=llvm::RTLIB::MUL_I64, Call_I128=llvm::RTLIB::MUL_I128)
> at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:2035
> #7 0x00007ffff21f55bd in (anonymous namespace)::SelectionDAGLegalize::ExpandNode (this=this at entry=0x7fffffffd4e0, Node=Node at entry=0xf35db0)
> at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:3495
> #8 0x00007ffff2200d86 in (anonymous namespace)::SelectionDAGLegalize::LegalizeOp (this=this at entry=0x7fffffffd4e0, Node=Node at entry=0xf35db0)
> at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:1364
> #9 0x00007ffff22026f9 in LegalizeDAG (this=0x7fffffffd4e0) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:237
> #10 llvm::SelectionDAG::Legalize (this=<optimized out>) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp:4074
> #11 0x00007ffff22e22c7 in llvm::SelectionDAGISel::CodeGenAndEmitDAG (this=this at entry=0xf89570) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:724
> #12 0x00007ffff22e2f37 in llvm::SelectionDAGISel::SelectBasicBlock (this=this at entry=0xf89570, Begin=..., Begin at entry=..., End=..., End at entry=..., HadTailCall=@0x7fffffffd740: false)
> at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:588
> #13 0x00007ffff22e6caf in llvm::SelectionDAGISel::SelectAllBasicBlocks (this=this at entry=0xf89570, Fn=...) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:1193
> #14 0x00007ffff22e8246 in llvm::SelectionDAGISel::runOnMachineFunction (this=0xf89570, mf=...) at /home/daenzer/src/llvm-git/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp:432
> #15 0x00007ffff1db8b0f in llvm::FPPassManager::runOnFunction (this=0x67bc00, F=...) at /home/daenzer/src/llvm-git/llvm/lib/IR/LegacyPassManager.cpp:1537
> #16 0x00007ffff1db8f4b in llvm::FPPassManager::runOnModule (this=0x67bc00, M=...) at /home/daenzer/src/llvm-git/llvm/lib/IR/LegacyPassManager.cpp:1557
> #17 0x00007ffff1db927d in runOnModule (M=..., this=<optimized out>) at /home/daenzer/src/llvm-git/llvm/lib/IR/LegacyPassManager.cpp:1615
> #18 llvm::legacy::PassManagerImpl::run (this=0xb86440, M=...) at /home/daenzer/src/llvm-git/llvm/lib/IR/LegacyPassManager.cpp:1710
> #19 0x00007ffff1db9459 in llvm::legacy::PassManager::run (this=this at entry=0x7fffffffda20, M=...) at /home/daenzer/src/llvm-git/llvm/lib/IR/LegacyPassManager.cpp:1745
> #20 0x00007ffff23782c6 in LLVMTargetMachineEmit (T=T at entry=0xf341d0, M=M at entry=0x6f7280, OS=..., codegen=codegen at entry=LLVMObjectFile, ErrorMessage=ErrorMessage at entry=0x7fffffffdb90)
> at /home/daenzer/src/llvm-git/llvm/lib/Target/TargetMachineC.cpp:232
> #21 0x00007ffff2378640 in LLVMTargetMachineEmitToMemoryBuffer (T=T at entry=0xf341d0, M=M at entry=0x6f7280, codegen=codegen at entry=LLVMObjectFile, ErrorMessage=ErrorMessage at entry=0x7fffffffdb90,
> OutMemBuf=OutMemBuf at entry=0x7fffffffdba0) at /home/daenzer/src/llvm-git/llvm/lib/Target/TargetMachineC.cpp:258
> #22 0x00007ffff08c80c8 in radeon_llvm_compile (M=M at entry=0x6f7280, binary=binary at entry=0x7fffffffdc80, gpu_family=0x7ffff0950dea "kaveri", dump=dump at entry=0)
> at ../../../../../src/gallium/drivers/radeon/radeon_llvm_emit.c:124
> #23 0x00007ffff08b5db9 in si_compile_llvm (sctx=sctx at entry=0x67fe90, shader=0x8fa310, mod=mod at entry=0x6f7280) at ../../../../../src/gallium/drivers/radeonsi/si_shader.c:2286
> #24 0x00007ffff08ae014 in si_create_compute_state (ctx=0x67fe90, cso=<optimized out>) at ../../../../../src/gallium/drivers/radeonsi/si_compute.c:79
> #25 0x00007ffff32d94af in clover::kernel::exec_context::bind (this=this at entry=0xf3b388, _q=0x0, _q at entry=0x67f9c0) at ../../../../../src/gallium/state_trackers/clover/core/kernel.cpp:189
> #26 0x00007ffff32d9fbc in clover::kernel::launch (this=0xf3b350, q=..., grid_offset=..., grid_size=..., block_size=...) at ../../../../../src/gallium/state_trackers/clover/core/kernel.cpp:72
> #27 0x00007ffff32d4f25 in clover::event::trigger (this=this at entry=0x9d4110) at ../../../../../src/gallium/state_trackers/clover/core/event.cpp:42
> #28 0x00007ffff32d59a8 in trigger (this=0x9d4110) at ../../../../../src/gallium/state_trackers/clover/core/event.cpp:85
> #29 clover::hard_event::hard_event(clover::command_queue&, unsigned int, clover::ref_vector<clover::event> const&, std::function<void (clover::event&)>) (this=0x9d4110, q=...,
> command=<optimized out>, deps=..., action=...) at ../../../../../src/gallium/state_trackers/clover/core/event.cpp:84
> #30 0x00007ffff32f4b38 in clEnqueueNDRangeKernel (d_q=0x67f9c8, d_kern=<optimized out>, dims=<optimized out>, d_grid_offset=<optimized out>, d_grid_size=<optimized out>,
> d_block_size=<optimized out>, num_deps=0, d_deps=0x0, rd_ev=0x0) at ../../../../../src/gallium/state_trackers/clover/api/kernel.cpp:283
> #31 0x00007ffff7bd1928 in piglit_cl_enqueue_ND_range_kernel (command_queue=0x67f9c8, kernel=0xf3b358, work_dim=1, global_work_size=0x7fffffffe2f8, local_work_size=0x7fffffffe310)
> at /home/daenzer/src/piglit-git/piglit/tests/util/piglit-util-cl.c:1068
> #32 0x00007ffff7bd19b3 in piglit_cl_execute_ND_range_kernel (command_queue=0x67f9c8, kernel=0xf3b358, work_dim=1, global_work_size=0x7fffffffe2f8, local_work_size=0x7fffffffe310)
> at /home/daenzer/src/piglit-git/piglit/tests/util/piglit-util-cl.c:1089
> #33 0x0000000000408081 in test_kernel (config=0x60cce0 <config>, env=0x7fffffffe400, test=...) at /home/daenzer/src/piglit-git/piglit/tests/cl/program/program-tester.c:1861
> #34 0x000000000040855d in piglit_cl_test (argc=2, argv=0x7fffffffe6b8, config=0x60cce0 <config>, env=0x7fffffffe400)
> at /home/daenzer/src/piglit-git/piglit/tests/cl/program/program-tester.c:1975
> #35 0x00007ffff7bd57f7 in piglit_cl_program_test_run (argc=2, argv=0x7fffffffe6b8, void_config=0x60cce0 <config>, version=11,
> platform_id=0x7ffff470fd00 <(anonymous namespace)::_clover_platform>, device_id=0x61e6d8) at /home/daenzer/src/piglit-git/piglit/tests/util/piglit-framework-cl-program.c:344
> #36 0x00007ffff7bd3e84 in piglit_cl_framework_run (argc=2, argv=0x7fffffffe6b8) at /home/daenzer/src/piglit-git/piglit/tests/util/piglit-framework-cl.c:338
> #37 0x0000000000402ff6 in main (argc=2, argv=0x7fffffffe6b8) at /home/daenzer/src/piglit-git/piglit/tests/cl/program/program-tester.c:162
>
> ; ModuleID = 'radeon'
> target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:64:64-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64"
> target triple = "r600--"
>
> ; Function Attrs: nounwind readonly
> declare float @llvm.pow.f32(float, float) #0
>
> ; Function Attrs: nounwind
> define void @rgb_gamma_u8_to_ragabaf(i8 addrspace(1)* nocapture readonly %in, <4 x float> addrspace(1)* nocapture %out) #1 {
> entry:
> %x.i.i = tail call i32 @llvm.r600.read.tgid.x() #2
> %x.i12.i = tail call i32 @llvm.r600.read.local.size.x() #2
> %mul26.i = mul i32 %x.i12.i, %x.i.i
> %x.i4.i = tail call i32 @llvm.r600.read.tidig.x() #2
> %add.i = add i32 %x.i4.i, %mul26.i
> %mul = mul nsw i32 %add.i, 3
> %0 = sext i32 %mul to i64
> %arrayidx = getelementptr inbounds i8 addrspace(1)* %in, i64 %0
> %1 = load i8 addrspace(1)* %arrayidx, align 1, !tbaa !3
> %add = add nsw i32 %mul, 1
> %2 = sext i32 %add to i64
> %arrayidx2 = getelementptr inbounds i8 addrspace(1)* %in, i64 %2
> %3 = load i8 addrspace(1)* %arrayidx2, align 1, !tbaa !3
> %add4 = add nsw i32 %mul, 2
> %4 = sext i32 %add4 to i64
> %arrayidx5 = getelementptr inbounds i8 addrspace(1)* %in, i64 %4
> %5 = load i8 addrspace(1)* %arrayidx5, align 1, !tbaa !3
> %conv.i.i.i = uitofp i8 %1 to float
> %6 = insertelement <3 x float> undef, float %conv.i.i.i, i32 0
> %conv.i4.i.i = uitofp i8 %3 to float
> %7 = insertelement <3 x float> %6, float %conv.i4.i.i, i32 1
> %conv.i.i = uitofp i8 %5 to float
> %vecinit2.i = insertelement <3 x float> %7, float %conv.i.i, i32 2
> %div = fdiv <3 x float> %vecinit2.i, <float 2.550000e+02, float 2.550000e+02, float 2.550000e+02>, !fpmath !6
> %8 = extractelement <3 x float> %div, i32 0
> %cmp.i50 = fcmp ogt float %8, 0x3FA41C8220000000
> br i1 %cmp.i50, label %if.then.i54, label %if.end.i56
>
> if.then.i54: ; preds = %entry
> %add.i51 = fadd float %8, 0x3FAC28F5C0000000
> %div.i52 = fdiv float %add.i51, 0x3FF0E147A0000000, !fpmath !6
> %call.i53 = tail call float @llvm.pow.f32(float %div.i52, float 0x4003333340000000) #3
> br label %gamma_2_2_to_linear.exit58
>
> if.end.i56: ; preds = %entry
> %div1.i55 = fdiv float %8, 0x4029D70A40000000, !fpmath !6
> br label %gamma_2_2_to_linear.exit58
>
> gamma_2_2_to_linear.exit58: ; preds = %if.end.i56, %if.then.i54
> %retval.0.i57 = phi float [ %call.i53, %if.then.i54 ], [ %div1.i55, %if.end.i56 ]
> %vecinit = insertelement <4 x float> undef, float %retval.0.i57, i32 0
> %9 = extractelement <3 x float> %div, i32 1
> %cmp.i41 = fcmp ogt float %9, 0x3FA41C8220000000
> br i1 %cmp.i41, label %if.then.i45, label %if.end.i47
>
> if.then.i45: ; preds = %gamma_2_2_to_linear.exit58
> %add.i42 = fadd float %9, 0x3FAC28F5C0000000
> %div.i43 = fdiv float %add.i42, 0x3FF0E147A0000000, !fpmath !6
> %call.i44 = tail call float @llvm.pow.f32(float %div.i43, float 0x4003333340000000) #3
> br label %gamma_2_2_to_linear.exit49
>
> if.end.i47: ; preds = %gamma_2_2_to_linear.exit58
> %div1.i46 = fdiv float %9, 0x4029D70A40000000, !fpmath !6
> br label %gamma_2_2_to_linear.exit49
>
> gamma_2_2_to_linear.exit49: ; preds = %if.end.i47, %if.then.i45
> %retval.0.i48 = phi float [ %call.i44, %if.then.i45 ], [ %div1.i46, %if.end.i47 ]
> %vecinit10 = insertelement <4 x float> %vecinit, float %retval.0.i48, i32 1
> %10 = extractelement <3 x float> %div, i32 2
> %cmp.i = fcmp ogt float %10, 0x3FA41C8220000000
> br i1 %cmp.i, label %if.then.i, label %if.end.i
>
> if.then.i: ; preds = %gamma_2_2_to_linear.exit49
> %add.i40 = fadd float %10, 0x3FAC28F5C0000000
> %div.i = fdiv float %add.i40, 0x3FF0E147A0000000, !fpmath !6
> %call.i = tail call float @llvm.pow.f32(float %div.i, float 0x4003333340000000) #3
> br label %gamma_2_2_to_linear.exit
>
> if.end.i: ; preds = %gamma_2_2_to_linear.exit49
> %div1.i = fdiv float %10, 0x4029D70A40000000, !fpmath !6
> br label %gamma_2_2_to_linear.exit
>
> gamma_2_2_to_linear.exit: ; preds = %if.end.i, %if.then.i
> %retval.0.i = phi float [ %call.i, %if.then.i ], [ %div1.i, %if.end.i ]
> %vecinit12 = insertelement <4 x float> %vecinit10, float %retval.0.i, i32 2
> %11 = insertelement <4 x float> %vecinit12, float 1.000000e+00, i32 3
> %12 = sext i32 %add.i to i64
> %arrayidx15 = getelementptr inbounds <4 x float> addrspace(1)* %out, i64 %12
> store <4 x float> %11, <4 x float> addrspace(1)* %arrayidx15, align 16, !tbaa !3
> ret void
> }
>
> ; Function Attrs: nounwind readnone
> declare i32 @llvm.r600.read.tgid.x() #2
>
> ; Function Attrs: nounwind readnone
> declare i32 @llvm.r600.read.local.size.x() #2
>
> ; Function Attrs: nounwind readnone
> declare i32 @llvm.r600.read.tidig.x() #2
>
> attributes #0 = { nounwind readonly }
> attributes #1 = { nounwind "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
> attributes #2 = { nounwind readnone }
> attributes #3 = { nounwind }
>
> !opencl.kernels = !{!0}
> !llvm.ident = !{!1, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2}
>
> !0 = metadata !{void (i8 addrspace(1)*, <4 x float> addrspace(1)*)* @rgb_gamma_u8_to_ragabaf}
> !1 = metadata !{metadata !"clang version 3.5 (http://llvm.org/git/clang.git 854ec99dfa5c443dd45451a73308095539f3ef4a) (llvm/trunk 201279)"}
> !2 = metadata !{metadata !"clang version 3.5 (http://llvm.org/git/clang.git 51910c637bb73218607144eef25e144bb9469342) (llvm/trunk 198901)"}
> !3 = metadata !{metadata !4, metadata !4, i64 0}
> !4 = metadata !{metadata !"omnipotent char", metadata !5, i64 0}
> !5 = metadata !{metadata !"Simple C/C++ TBAA"}
> !6 = metadata !{float 2.500000e+00}
> 0x199c4f0: i32 = ExternalSymbol'__muldi3'
> Undefined function
> UNREACHABLE executed at /home/daenzer/src/llvm-git/llvm/lib/Target/R600/AMDGPUISelLowering.h:96!
> _______________________________________________
> llvm-commits mailing list
> llvm-commits at cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits
More information about the llvm-commits
mailing list