Author: jlebar Date: Mon Sep 25 12:41:56 2017 New Revision: 314142 URL: http://llvm.org/viewvc/llvm-project?rev=314142&view=rev Log: Revert "[NVPTX] added match.{any,all}.sync instructions, intrinsics & builtins.", rL314135.
Causing assertion failures on macos: > Assertion failed: (Num < NumOperands && "Invalid child # of SDNode!"), > function getOperand, file > /Users/buildslave/jenkins/workspace/clang-stage1-cmake-RA-incremental/llvm/include/llvm/CodeGen/SelectionDAGNodes.h, > line 835. http://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/42739/testReport/LLVM/CodeGen_NVPTX/surf_read_cuda_ll/ Modified: cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def cfe/trunk/lib/CodeGen/CGBuiltin.cpp cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h cfe/trunk/test/CodeGen/builtins-nvptx-ptx60.cu Modified: cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def?rev=314142&r1=314141&r2=314142&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsNVPTX.def Mon Sep 25 12:41:56 2017 @@ -413,13 +413,6 @@ TARGET_BUILTIN(__nvvm_vote_any_sync, "bU TARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib", "", "ptx60") TARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib", "", "ptx60") -// Match -TARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi", "", "ptx60") -TARGET_BUILTIN(__nvvm_match_any_sync_i64, "WiUiWi", "", "ptx60") -// These return a pair {value, predicate}, which requires custom lowering. -TARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*", "", "ptx60") -TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", "ptx60") - // Membar BUILTIN(__nvvm_membar_cta, "v", "") Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=314142&r1=314141&r2=314142&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original) +++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Mon Sep 25 12:41:56 2017 @@ -9589,21 +9589,6 @@ Value *CodeGenFunction::EmitNVPTXBuiltin {Ptr->getType()->getPointerElementType(), Ptr->getType()}), {Ptr, EmitScalarExpr(E->getArg(1)), EmitScalarExpr(E->getArg(2))}); } - case NVPTX::BI__nvvm_match_all_sync_i32p: - case NVPTX::BI__nvvm_match_all_sync_i64p: { - Value *Mask = EmitScalarExpr(E->getArg(0)); - Value *Val = EmitScalarExpr(E->getArg(1)); - Address PredOutPtr = EmitPointerWithAlignment(E->getArg(2)); - Value *ResultPair = Builder.CreateCall( - CGM.getIntrinsic(BuiltinID == NVPTX::BI__nvvm_match_all_sync_i32p - ? Intrinsic::nvvm_match_all_sync_i32p - : Intrinsic::nvvm_match_all_sync_i64p), - {Mask, Val}); - Value *Pred = Builder.CreateZExt(Builder.CreateExtractValue(ResultPair, 1), - PredOutPtr.getElementType()); - Builder.CreateStore(Pred, PredOutPtr); - return Builder.CreateExtractValue(ResultPair, 0); - } default: return nullptr; } Modified: cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h?rev=314142&r1=314141&r2=314142&view=diff ============================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h Mon Sep 25 12:41:56 2017 @@ -92,9 +92,8 @@ __MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_ #endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300 -#if CUDA_VERSION >= 9000 -#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300) // __shfl_sync_* variants available in CUDA-9 +#if CUDA_VERSION >= 9000 && (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300) #pragma push_macro("__MAKE_SYNC_SHUFFLES") #define __MAKE_SYNC_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, \ __Mask) \ @@ -188,33 +187,8 @@ inline __device__ unsigned int __ballot_ inline __device__ unsigned int activemask() { return __nvvm_vote_ballot(1); } -#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300 - -// Define __match* builtins CUDA-9 headers expect to see. -#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700 -inline __device__ unsigned int __match32_any_sync(unsigned int mask, - unsigned int value) { - return __nvvm_match_any_sync_i32(mask, value); -} - -inline __device__ unsigned long long -__match64_any_sync(unsigned int mask, unsigned long long value) { - return __nvvm_match_any_sync_i64(mask, value); -} - -inline __device__ unsigned int -__match32_all_sync(unsigned int mask, unsigned int value, int *pred) { - return __nvvm_match_all_sync_i32p(mask, value, pred); -} - -inline __device__ unsigned long long -__match64_all_sync(unsigned int mask, unsigned long long value, int *pred) { - return __nvvm_match_all_sync_i64p(mask, value, pred); -} -#include "crt/sm_70_rt.hpp" - -#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700 -#endif // __CUDA_VERSION >= 9000 +#endif // __CUDA_VERSION >= 9000 && (!defined(__CUDA_ARCH__) || + // __CUDA_ARCH__ >= 300) // sm_32 intrinsics: __ldg and __funnelshift_{l,lc,r,rc}. Modified: cfe/trunk/test/CodeGen/builtins-nvptx-ptx60.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/builtins-nvptx-ptx60.cu?rev=314142&r1=314141&r2=314142&view=diff ============================================================================== --- cfe/trunk/test/CodeGen/builtins-nvptx-ptx60.cu (original) +++ cfe/trunk/test/CodeGen/builtins-nvptx-ptx60.cu Mon Sep 25 12:41:56 2017 @@ -10,8 +10,6 @@ #define __shared__ __attribute__((shared)) #define __constant__ __attribute__((constant)) -typedef unsigned long long uint64_t; - // We have to keep all builtins that depend on particular target feature in the // same function, because the codegen will stop after the very first function // that encounters an error, so -verify will not be able to find errors in @@ -19,8 +17,7 @@ typedef unsigned long long uint64_t; // CHECK-LABEL: nvvm_sync __device__ void nvvm_sync(unsigned mask, int i, float f, int a, int b, - bool pred, uint64_t i64) { - + bool pred) { // CHECK: call void @llvm.nvvm.bar.warp.sync(i32 // expected-error@+1 {{'__nvvm_bar_warp_sync' needs target feature ptx60}} __nvvm_bar_warp_sync(mask); @@ -76,22 +73,5 @@ __device__ void nvvm_sync(unsigned mask, // expected-error@+1 {{'__nvvm_vote_ballot_sync' needs target feature ptx60}} __nvvm_vote_ballot_sync(mask, pred); - // - // MATCH.{ALL,ANY}.SYNC - // - - // CHECK: call i32 @llvm.nvvm.match.any.sync.i32(i32 - // expected-error@+1 {{'__nvvm_match_any_sync_i32' needs target feature ptx60}} - __nvvm_match_any_sync_i32(mask, i); - // CHECK: call i64 @llvm.nvvm.match.any.sync.i64(i32 - // expected-error@+1 {{'__nvvm_match_any_sync_i64' needs target feature ptx60}} - __nvvm_match_any_sync_i64(mask, i64); - // CHECK: call { i32, i1 } @llvm.nvvm.match.all.sync.i32p(i32 - // expected-error@+1 {{'__nvvm_match_all_sync_i32p' needs target feature ptx60}} - __nvvm_match_all_sync_i32p(mask, i, &i); - // CHECK: call { i64, i1 } @llvm.nvvm.match.all.sync.i64p(i32 - // expected-error@+1 {{'__nvvm_match_all_sync_i64p' needs target feature ptx60}} - __nvvm_match_all_sync_i64p(mask, i64, &i); - // CHECK: ret void } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
