-
Notifications
You must be signed in to change notification settings - Fork 1.6k
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Build Error: "undefined reference to `pthread_rwlock_{destroy,rdlock,init,unlock,wrlock}' #14
Comments
Oh. I think that clBLAS has been adding and removing pthread as a dependency a couple of times. Your version of clBlas seems to require pthread while mine doesn't. ISAAC itself doesn't depend on pthread. I've modified the build system so that it always links against pthread when clBlas is detected. Can you pull and retry? |
Wow, fast turnaround, thank you! I'll pull and retry within the hour. :)
…On 20 January 2017 22:23:03 GMT+00:00, Phillippe Tillet ***@***.***> wrote:
Oh. I think that clBLAS has been adding and removing pthread as a
dependency a couple of time. Your version of clBlas that cmake wants to
benchmark again seems to require pthread. Isaac itself doesn't depend
on pthread.
I've modified the build system so that it always links against pthread
when clBlas is detected. Can you pull and retry?
--
You are receiving this because you authored the thread.
Reply to this email directly or view it on GitHub:
#14 (comment)
--
Sent from my Android device with K-9 Mail. Please excuse my brevity.
|
It built! :) Now, I have a new problem.. if I immediately follow that up with Perhaps this merits another issue.. |
Is the issue solved with the segfault? This usually comes from a machine configuration issue (such as launching on machines with bumblebee without optirun) |
…rization fix test_vectorization and test_load_cache_modifier
* Update * Update * Update * Update * Update * Update
fix test_vectorization and test_load_cache_modifier
When running [convert_blocked1d_to_slice0](https://github.com/triton-lang/triton/blob/0ba5f0c3cd029d5c3d1f01b9bf29dac32c27345e/test/Conversion/tritongpu_to_llvm.mlir#L924) Triton ends up computing a rank of a matrix with 0 columns during linear layout lowering, which trips up f2reduce, and causes undefined behavior, detectable through [UBSAN](https://clang.llvm.org/docs/UndefinedBehaviorSanitizer.html). Fix this by returning the rank (0) early in these cases, without calling f2reduce. <details><summary>Stack trace</summary> <p> ``` third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30: runtime error: shift exponent 18446744073709551615 is too large for 64-bit type 'unsigned long long' #0 0x556ee2fea3be in inplace_rref_small third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30 #1 0x556ee2fea3be in f2reduce::inplace_rref_strided(unsigned long*, unsigned long, unsigned long, unsigned long) third_party/triton/third_party/f2reduce/f2reduce.cpp:470:9 #2 0x556ee2ea70da in getMatrixRank third_party/triton/lib/Tools/LinearLayout.cpp:125:3 #3 0x556ee2ea70da in mlir::triton::LinearLayout::checkInvariants(bool) third_party/triton/lib/Tools/LinearLayout.cpp:299:7 #4 0x556ee2ea656d in mlir::triton::LinearLayout::tryCreate(llvm::MapVector<mlir::StringAttr, std::__u::vector<std::__u::vector<int, std::__u::allocator<int>>, std::__u::allocator<std::__u::vector<int, std::__u::allocator<int>>>>, llvm::DenseMap<mlir::StringAttr, unsigned int, llvm::DenseMapInfo<mlir::StringAttr, void>, llvm::detail::DenseMapPair<mlir::StringAttr, unsigned int>>, llvm::SmallVector<std::__u::pair<mlir::StringAttr, std::__u::vector<std::__u::vector<int, std::__u::allocator<int>>, std::__u::allocator<std::__u::vector<int, std::__u::allocator<int>>>>>, 0u>>, llvm::ArrayRef<std::__u::pair<mlir::StringAttr, int>>, bool) third_party/triton/lib/Tools/LinearLayout.cpp:190:41 #5 0x556ee2eb2150 in mlir::triton::LinearLayout::divideRight(mlir::triton::LinearLayout const&) third_party/triton/lib/Tools/LinearLayout.cpp:654:51 #6 0x556ee2ee1c39 in mlir::cvtNeedsSharedMemory(mlir::RankedTensorType, mlir::RankedTensorType) third_party/triton/lib/Analysis/Utility.cpp:652:14 #7 0x556ee2cf38fd in mlir::triton::getRepShapeForCvtLayout(mlir::triton::gpu::ConvertLayoutOp) third_party/triton/lib/Analysis/Allocation.cpp:66:8 #8 0x556ee2cf3efa in mlir::triton::getScratchConfigForCvtLayout(mlir::triton::gpu::ConvertLayoutOp, unsigned int&, unsigned int&) third_party/triton/lib/Analysis/Allocation.cpp:95:19 #9 0x556ee2cf6057 in mlir::triton::AllocationAnalysis::getScratchValueSize(mlir::Operation*) third_party/triton/lib/Analysis/Allocation.cpp:272:24 #10 0x556ee2cf5499 in operator() third_party/triton/lib/Analysis/Allocation.cpp:343:7 #11 0x556ee2cf5499 in void llvm::function_ref<void (mlir::Operation*)>::callback_fn<mlir::triton::AllocationAnalysis::getValuesAndSizes()::'lambda'(mlir::Operation*)>(long, mlir::Operation*) third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12 #12 0x556edeeee7a9 in operator() third_party/llvm/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12 #13 0x556edeeee7a9 in void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:174:5 #14 0x556edeeee87c in void mlir::detail::walk<mlir::ForwardIterator>(mlir::Operation*, llvm::function_ref<void (mlir::Operation*)>, mlir::WalkOrder) third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:182:9 #15 0x556ee2cf49e7 in walk<(mlir::WalkOrder)0, mlir::ForwardIterator, (lambda at third_party/triton/lib/Analysis/Allocation.cpp:341:42), mlir::Operation *, void> third_party/llvm/llvm-project/mlir/include/mlir/IR/Visitors.h:313:10 #16 0x556ee2cf49e7 in walk<(mlir::WalkOrder)0, mlir::ForwardIterator, (lambda at third_party/triton/lib/Analysis/Allocation.cpp:341:42), void> third_party/llvm/llvm-project/mlir/include/mlir/IR/Operation.h:794:12 #17 0x556ee2cf49e7 in mlir::triton::AllocationAnalysis::getValuesAndSizes() third_party/triton/lib/Analysis/Allocation.cpp:341:16 #18 0x556ee2cf4852 in run third_party/triton/lib/Analysis/Allocation.cpp:182:5 #19 0x556ee2cf4852 in AllocationAnalysis third_party/triton/lib/Analysis/Allocation.cpp:169:5 #20 0x556ee2cf4852 in mlir::Allocation::run(llvm::DenseMap<mlir::FunctionOpInterface, mlir::Allocation, llvm::DenseMapInfo<mlir::FunctionOpInterface, void>, llvm::detail::DenseMapPair<mlir::FunctionOpInterface, mlir::Allocation>>&) third_party/triton/lib/Analysis/Allocation.cpp:627:3 #21 0x556ee1677402 in operator() third_party/triton/include/triton/Analysis/Allocation.h:227:26 #22 0x556ee1677402 in void mlir::CallGraph<mlir::Allocation>::doWalk<(mlir::WalkOrder)0, (mlir::WalkOrder)1, mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::CallOpInterface, mlir::FunctionOpInterface), mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::FunctionOpInterface)>(mlir::FunctionOpInterface, llvm::DenseSet<mlir::FunctionOpInterface, llvm::DenseMapInfo<mlir::FunctionOpInterface, void>>&, mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::CallOpInterface, mlir::FunctionOpInterface), mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp)::'lambda'(mlir::FunctionOpInterface)) third_party/triton/include/triton/Analysis/Utility.h:350:7 #23 0x556ee16756b3 in walk<(mlir::WalkOrder)0, (mlir::WalkOrder)1, (lambda at third_party/triton/include/triton/Analysis/Allocation.h:222:9), (lambda at third_party/triton/include/triton/Analysis/Allocation.h:224:9)> third_party/triton/include/triton/Analysis/Utility.h:242:7 #24 0x556ee16756b3 in mlir::ModuleAllocation::ModuleAllocation(mlir::ModuleOp) third_party/triton/include/triton/Analysis/Allocation.h:220:5 #25 0x556ee2c2bf18 in (anonymous namespace)::AllocateSharedMemory::runOnOperation() third_party/triton/lib/Conversion/TritonGPUToLLVM/AllocateSharedMemory.cpp:26:22 ... UndefinedBehaviorSanitizer: invalid-shift-exponent third_party/triton/third_party/f2reduce/f2reduce.cpp:421:30 ``` </p> </details>
This PR fixes triton-lang#1176 IGC detects the call of `__devicelib_assert_fail` and replace it with a 'safe' implementation. However, the SYCL library contains a 'fallback' implementation of assertion, which does not work in our setup. If we mark the function with `InternalLinkage`, the fallback implementation is inlined and IGC cannot replace it with the safe implementation. By declaring `__devicelib_assert_fail` as an external function in SYCL library, IGC can correctly insert its implementation. The diff between the old and new `libsycl-spir64-unknown-unknown.ll` is as follows: ```diff @@ -5424,149 +5424,7 @@ declare extern_weak dso_local spir_func noundef i32 @_Z18__spirv_AtomicLoadPU3AS declare void @llvm.memcpy.p4.p1.i64(ptr addrspace(4) noalias nocapture writeonly, ptr addrspace(1) noalias nocapture readonly, i64, i1 immarg) triton-lang#16 ; Function Attrs: convergent mustprogress norecurse nounwind -define weak dso_local spir_func void @__devicelib_assert_fail(ptr addrspace(4) noundef %0, ptr addrspace(4) noundef %1, i32 noundef %2, ptr addrspace(4) noundef %3, i64 noundef %4, i64 noundef %5, i64 noundef %6, i64 noundef %7, i64 noundef %8, i64 noundef %9) local_unnamed_addr triton-lang#14 !srcloc !720 { - %11 = tail call spir_func noundef i32 @_Z29__spirv_AtomicCompareExchangePU3AS1iN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_ii(ptr addrspace(1) noundef @SPIR_AssertHappenedMem, i32 noundef 1, i32 noundef 16, i32 noundef 16, i32 noundef 1, i32 noundef 0) triton-lang#54 - %12 = icmp eq i32 %11, 0 - br i1 %12, label %13, label %92 - -13: ; preds = %10 - store i32 %2, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 4), align 8, !tbaa !721 - store i64 %4, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 5), align 8, !tbaa !722 - store i64 %5, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 6), align 8, !tbaa !723 - store i64 %6, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 7), align 8, !tbaa !724 - store i64 %7, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 8), align 8, !tbaa !725 - store i64 %8, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 9), align 8, !tbaa !726 - store i64 %9, ptr addrspace(1) getelementptr inbounds (%struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 10), align 8, !tbaa !727 - %14 = icmp eq ptr addrspace(4) %0, null - br i1 %14, label %23, label %15 - -15: ; preds = %20, %13 - %16 = phi i32 [ %22, %20 ], [ 0, %13 ] - %17 = phi ptr addrspace(4) [ %21, %20 ], [ %0, %13 ] - %18 = load i8, ptr addrspace(4) %17, align 1, !tbaa !718 - %19 = icmp eq i8 %18, 0 - br i1 %19, label %23, label %20 - -20: ; preds = %15 - %21 = getelementptr inbounds i8, ptr addrspace(4) %17, i64 1 - %22 = add nuw nsw i32 %16, 1 - br label %15, !llvm.loop !728 - -23: ; preds = %15, %13 - %24 = phi i32 [ 0, %13 ], [ %16, %15 ] - %25 = icmp eq ptr addrspace(4) %1, null - br i1 %25, label %34, label %26 - -26: ; preds = %31, %23 - %27 = phi i32 [ %33, %31 ], [ 0, %23 ] - %28 = phi ptr addrspace(4) [ %32, %31 ], [ %1, %23 ] - %29 = load i8, ptr addrspace(4) %28, align 1, !tbaa !718 - %30 = icmp eq i8 %29, 0 - br i1 %30, label %34, label %31 - -31: ; preds = %26 - %32 = getelementptr inbounds i8, ptr addrspace(4) %28, i64 1 - %33 = add nuw nsw i32 %27, 1 - br label %26, !llvm.loop !729 - -34: ; preds = %26, %23 - %35 = phi i32 [ 0, %23 ], [ %27, %26 ] - %36 = icmp eq ptr addrspace(4) %3, null - br i1 %36, label %37, label %40 - -37: ; preds = %34 - %38 = tail call i32 @llvm.umin.i32(i32 %24, i32 256) - %39 = tail call i32 @llvm.umin.i32(i32 %35, i32 256) - br label %52 - -40: ; preds = %45, %34 - %41 = phi i32 [ %47, %45 ], [ 0, %34 ] - %42 = phi ptr addrspace(4) [ %46, %45 ], [ %3, %34 ] - %43 = load i8, ptr addrspace(4) %42, align 1, !tbaa !718 - %44 = icmp eq i8 %43, 0 - br i1 %44, label %48, label %45 - -45: ; preds = %40 - %46 = getelementptr inbounds i8, ptr addrspace(4) %42, i64 1 - %47 = add i32 %41, 1 - br label %40, !llvm.loop !730 - -48: ; preds = %40 - %49 = tail call i32 @llvm.umin.i32(i32 %24, i32 256) - %50 = tail call i32 @llvm.umin.i32(i32 %35, i32 256) - %51 = tail call i32 @llvm.umin.i32(i32 %41, i32 128) - br label %52 - -52: ; preds = %48, %37 - %53 = phi i32 [ %39, %37 ], [ %50, %48 ] - %54 = phi i32 [ %38, %37 ], [ %49, %48 ] - %55 = phi i32 [ 0, %37 ], [ %51, %48 ] - br label %56 - -56: ; preds = %62, %52 - %57 = phi i32 [ 0, %52 ], [ %67, %62 ] - %58 = icmp ult i32 %57, %54 - br i1 %58, label %62, label %59 - -59: ; preds = %56 - %60 = zext nneg i32 %54 to i64 - %61 = getelementptr inbounds %struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 1, i64 %60 - store i8 0, ptr addrspace(1) %61, align 1, !tbaa !718 - br label %68 - -62: ; preds = %56 - %63 = sext i32 %57 to i64 - %64 = getelementptr inbounds i8, ptr addrspace(4) %0, i64 %63 - %65 = load i8, ptr addrspace(4) %64, align 1, !tbaa !718 - %66 = getelementptr inbounds %struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 1, i64 %63 - store i8 %65, ptr addrspace(1) %66, align 1, !tbaa !718 - %67 = add nuw nsw i32 %57, 1 - br label %56, !llvm.loop !731 - -68: ; preds = %74, %59 - %69 = phi i32 [ 0, %59 ], [ %79, %74 ] - %70 = icmp ult i32 %69, %53 - br i1 %70, label %74, label %71 - -71: ; preds = %68 - %72 = zext nneg i32 %53 to i64 - %73 = getelementptr inbounds %struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 2, i64 %72 - store i8 0, ptr addrspace(1) %73, align 1, !tbaa !718 - br label %80 - -74: ; preds = %68 - %75 = sext i32 %69 to i64 - %76 = getelementptr inbounds i8, ptr addrspace(4) %1, i64 %75 - %77 = load i8, ptr addrspace(4) %76, align 1, !tbaa !718 - %78 = getelementptr inbounds %struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 2, i64 %75 - store i8 %77, ptr addrspace(1) %78, align 1, !tbaa !718 - %79 = add nuw nsw i32 %69, 1 - br label %68, !llvm.loop !732 - -80: ; preds = %86, %71 - %81 = phi i32 [ 0, %71 ], [ %91, %86 ] - %82 = icmp ult i32 %81, %55 - br i1 %82, label %86, label %83 - -83: ; preds = %80 - %84 = sext i32 %55 to i64 - %85 = getelementptr inbounds %struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 3, i64 %84 - store i8 0, ptr addrspace(1) %85, align 1, !tbaa !718 - tail call spir_func void @_Z19__spirv_AtomicStorePU3AS1iN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEi(ptr addrspace(1) noundef @SPIR_AssertHappenedMem, i32 noundef 1, i32 noundef 16, i32 noundef 2) triton-lang#54 - br label %92 - -86: ; preds = %80 - %87 = sext i32 %81 to i64 - %88 = getelementptr inbounds i8, ptr addrspace(4) %3, i64 %87 - %89 = load i8, ptr addrspace(4) %88, align 1, !tbaa !718 - %90 = getelementptr inbounds %struct.AssertHappened, ptr addrspace(1) @SPIR_AssertHappenedMem, i64 0, i32 3, i64 %87 - store i8 %89, ptr addrspace(1) %90, align 1, !tbaa !718 - %91 = add nuw nsw i32 %81, 1 - br label %80, !llvm.loop !733 - -92: ; preds = %83, %10 - ret void -} +declare extern_weak dso_local spir_func void @__devicelib_assert_fail(ptr addrspace(4) noundef %0, ptr addrspace(4) noundef %1, i32 noundef %2, ptr addrspace(4) noundef %3, i64 noundef %4, i64 noundef %5, i64 noundef %6, i64 noundef %7, i64 noundef %8, i64 noundef %9) local_unnamed_addr triton-lang#14 ; Function Attrs: convergent nounwind declare extern_weak dso_local spir_func noundef i32 @_Z29__spirv_AtomicCompareExchangePU3AS1iN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagES5_ii(ptr addrspace(1) noundef, i32 noundef, i32 noundef, i32 noundef, i32 noundef, i32 noundef) local_unnamed_addr triton-lang#15 ```
Signed-off-by: Ilya Enkovich <ilya.enkovich@intel.com>
Simple loop test based on Thomas' dot test
Signed-off-by: Ilya Enkovich <ilya.enkovich@intel.com>
I follow the build instructions on the Readme:
..and I get this build error:
I don't see a reference in a non-binary file to 'pthread' when I
sift
for it, but it's matched in lots of object files. So, perhaps it's just missing somewhere from the Makefiles...but I don't grok the C build system. :/The text was updated successfully, but these errors were encountered: