[TF] [XLA] Fix a number of compiler warnings on Mac OS X.
* Mostly these are unused/dead code. * Fix an ignored Status in shape_util.h PiperOrigin-RevId: 222126279
This commit is contained in:
parent
ea3ee9804d
commit
ee20d8c029
@ -104,7 +104,7 @@ std::unique_ptr<Array2D<NativeT>> MakeLinspaceArray2D(double from, double to,
|
||||
int64 count = n1 * n2;
|
||||
NativeT step =
|
||||
static_cast<NativeT>((count > 1) ? (to - from) / (count - 1) : 0);
|
||||
auto set = [&array, n1, n2](int64 index, NativeT value) {
|
||||
auto set = [&array, n2](int64 index, NativeT value) {
|
||||
(*array)(index / n2, index % n2) = value;
|
||||
};
|
||||
for (int64 i = 0; i < count - 1; ++i) {
|
||||
|
@ -746,8 +746,7 @@ StatusOr<std::unique_ptr<BufferAssignment>> BufferAssigner::Run(
|
||||
LogicalBuffer::AlignmentFunction color_alignment,
|
||||
bool allow_input_output_aliasing, bool allocate_buffers_for_constants,
|
||||
BufferLiveness::Colorer colorer, ReuseAllocationFunction reuse_checker) {
|
||||
BufferAssigner assigner(allow_input_output_aliasing,
|
||||
allocate_buffers_for_constants, std::move(colorer),
|
||||
BufferAssigner assigner(allocate_buffers_for_constants, std::move(colorer),
|
||||
std::move(reuse_checker));
|
||||
return assigner.CreateAssignment(module, std::move(hlo_ordering),
|
||||
std::move(buffer_size),
|
||||
|
@ -545,12 +545,10 @@ class BufferAssigner {
|
||||
ReuseAllocationFunction reuse_checker = nullptr);
|
||||
|
||||
private:
|
||||
BufferAssigner(bool allow_input_output_aliasing,
|
||||
bool allocate_buffers_for_constants,
|
||||
BufferAssigner(bool allocate_buffers_for_constants,
|
||||
BufferLiveness::Colorer colorer,
|
||||
ReuseAllocationFunction reuse_checker)
|
||||
: allow_input_output_aliasing_(allow_input_output_aliasing),
|
||||
allocate_buffers_for_constants_(allocate_buffers_for_constants),
|
||||
: allocate_buffers_for_constants_(allocate_buffers_for_constants),
|
||||
colorer_(colorer),
|
||||
reuse_checker_(reuse_checker) {}
|
||||
virtual ~BufferAssigner() = default;
|
||||
@ -640,10 +638,6 @@ class BufferAssigner {
|
||||
LogicalBuffer::Color::Hasher>
|
||||
SplitBuffersByColor(const absl::flat_hash_set<const LogicalBuffer*>& buffers);
|
||||
|
||||
// If true, buffer assignments assumes that input parameter buffers and output
|
||||
// buffers can be shared if their sizes match.
|
||||
bool allow_input_output_aliasing_;
|
||||
|
||||
// If true, allocate buffers for constant instructions.
|
||||
bool allocate_buffers_for_constants_;
|
||||
|
||||
|
@ -105,8 +105,6 @@ class ComputationPlacer {
|
||||
// Map from platform kind to computation placer singleton.
|
||||
static std::map<se::Platform::Id, State>* GetPlatformComputationPlacers();
|
||||
|
||||
se::Platform::Id platform_id_;
|
||||
|
||||
TF_DISALLOW_COPY_AND_ASSIGN(ComputationPlacer);
|
||||
};
|
||||
|
||||
|
@ -442,7 +442,6 @@ class CopyRemover {
|
||||
const HloOrdering& ordering, HloModule* module)
|
||||
: module_(module),
|
||||
alias_analysis_(alias_analysis),
|
||||
ordering_(ordering),
|
||||
buffer_value_tracker_(*module, alias_analysis, ordering) {}
|
||||
|
||||
// Try to elide the given copy. The copy is elided if the instruction is not
|
||||
@ -1003,7 +1002,6 @@ class CopyRemover {
|
||||
|
||||
HloModule* module_;
|
||||
const HloAliasAnalysis& alias_analysis_;
|
||||
const HloOrdering& ordering_;
|
||||
|
||||
// Object tracking the HLO values contained in each HLO buffer.
|
||||
BufferValueTracker buffer_value_tracker_;
|
||||
|
@ -61,17 +61,15 @@ Disabling these as a starting point.
|
||||
// TODO(b/64227304) Creating a custom pass pipeline will replace this.
|
||||
|
||||
namespace {
|
||||
|
||||
// TODO(sanjoy): remove this class.
|
||||
class FilteredFunctionPassManager : public llvm::legacy::FunctionPassManager {
|
||||
public:
|
||||
FilteredFunctionPassManager(llvm::Module* m, bool disable_expensive_passes)
|
||||
: llvm::legacy::FunctionPassManager(m),
|
||||
disable_expensive_passes_(disable_expensive_passes) {}
|
||||
explicit FilteredFunctionPassManager(llvm::Module* m)
|
||||
: llvm::legacy::FunctionPassManager(m) {}
|
||||
void add(llvm::Pass* p) override {
|
||||
llvm::legacy::FunctionPassManager::add(p);
|
||||
}
|
||||
|
||||
private:
|
||||
bool disable_expensive_passes_;
|
||||
};
|
||||
|
||||
class FilteredPassManager : public llvm::legacy::PassManager {
|
||||
@ -96,8 +94,7 @@ class FilteredPassManager : public llvm::legacy::PassManager {
|
||||
std::unique_ptr<llvm::MemoryBuffer> CompilerFunctor::operator()(
|
||||
llvm::Module& module) const {
|
||||
FilteredPassManager module_passes(disable_expensive_passes_);
|
||||
FilteredFunctionPassManager function_passes(&module,
|
||||
disable_expensive_passes_);
|
||||
FilteredFunctionPassManager function_passes(&module);
|
||||
|
||||
VLOG(2) << "IR before optimizations";
|
||||
XLA_VLOG_LINES(2, llvm_ir::DumpModuleToString(module));
|
||||
|
@ -22,7 +22,6 @@ limitations under the License.
|
||||
namespace {
|
||||
|
||||
const char* const kXlaOptimizeForSizeCpuOption = "xla_cpu_optimize_for_size";
|
||||
const char* const kXlaDisableVectorizedReduce = "xla_disable_vectorized_reduce";
|
||||
const char* const kLlvmIrDotTilingFactor = "xla_llvm_dot_tiling_factor";
|
||||
const char* const kXlaEnableExperimentalLlvmIrGemm =
|
||||
"xla_enable_experimental_llvm_ir_gemm";
|
||||
|
@ -140,7 +140,7 @@ StatusOr<llvm::Function*> IrEmitter::EmitComputation(
|
||||
// readcyclecounter if it is unavailable.
|
||||
bool use_rdtscp = arch_type_ == llvm::Triple::ArchType::x86 ||
|
||||
arch_type_ == llvm::Triple::ArchType::x86_64;
|
||||
profiling_state_ = ProfilingState(use_rdtscp, GetProfileCountersArgument());
|
||||
profiling_state_ = ProfilingState(use_rdtscp);
|
||||
if (instruction_order == nullptr) {
|
||||
TF_RETURN_IF_ERROR(computation->Accept(this));
|
||||
} else {
|
||||
@ -1379,33 +1379,6 @@ Status IrEmitter::HandleCrossReplicaSum(HloInstruction* crs) {
|
||||
return Status::OK();
|
||||
}
|
||||
|
||||
// Fills up the free variables in 'index_with_free_var' with values from
|
||||
// 'filler_index'. The size of free variables must be the same as the
|
||||
// size of 'filler_index'.
|
||||
//
|
||||
// This is often used after dimension reduction, where
|
||||
// 'index_with_free_var' has one or more dimensions reduced, which serves as
|
||||
// free variables (represented as nullptr). For example, if we have a 4
|
||||
// dimensional input and index for the dimension being reduced is
|
||||
// 2 (third dimension), we will have an index like [i, j, NULL, k]
|
||||
// after reduced dimension.
|
||||
//
|
||||
// Here we fill up that free variable by 'filler_index', which contains
|
||||
// the value in the reduced dimension.
|
||||
static llvm_ir::IrArray::Index FillReducedDimensionIndex(
|
||||
llvm_ir::IrArray::Index index_with_free_var,
|
||||
llvm_ir::IrArray::Index filler_index) {
|
||||
llvm_ir::IrArray::Index::const_iterator it = filler_index.begin();
|
||||
|
||||
for (size_t i = 0; i < index_with_free_var.size(); ++i) {
|
||||
if (index_with_free_var[i] == nullptr) {
|
||||
index_with_free_var[i] = *it++;
|
||||
}
|
||||
}
|
||||
CHECK(filler_index.end() == it);
|
||||
return index_with_free_var;
|
||||
}
|
||||
|
||||
Status IrEmitter::HandleParameter(HloInstruction* parameter) {
|
||||
VLOG(2) << "HandleParameter: " << parameter->ToString();
|
||||
return EmitTargetAddressForOp(parameter);
|
||||
@ -2194,14 +2167,6 @@ Status IrEmitter::HandlePad(HloInstruction* pad) {
|
||||
return Status::OK();
|
||||
}
|
||||
|
||||
// If `hlo` is a Transpose, returns its operand; otherwise returns `hlo` itself.
|
||||
static const HloInstruction* StripTranspose(const HloInstruction& hlo) {
|
||||
if (hlo.IsRank2Transpose()) {
|
||||
return hlo.operand(0);
|
||||
}
|
||||
return &hlo;
|
||||
}
|
||||
|
||||
Status IrEmitter::HandleFusion(HloInstruction* fusion) {
|
||||
auto* root = fusion->fused_expression_root();
|
||||
if (llvm_ir::CanEmitFusedDynamicUpdateSliceInPlace(fusion, assignment_)) {
|
||||
|
@ -467,9 +467,8 @@ class IrEmitter : public DfsHloVisitorWithDefault,
|
||||
// profiling a computation.
|
||||
class ProfilingState {
|
||||
public:
|
||||
ProfilingState() : use_rdtscp_(false), prof_counters_(nullptr) {}
|
||||
ProfilingState(bool use_rdtscp, llvm::Value* prof_counters)
|
||||
: use_rdtscp_(use_rdtscp), prof_counters_(prof_counters) {}
|
||||
ProfilingState() : use_rdtscp_(false) {}
|
||||
explicit ProfilingState(bool use_rdtscp) : use_rdtscp_(use_rdtscp) {}
|
||||
|
||||
// Record the cycle counter before an HLO executes.
|
||||
void RecordCycleStart(llvm::IRBuilder<>* b, HloInstruction* hlo);
|
||||
@ -494,9 +493,6 @@ class IrEmitter : public DfsHloVisitorWithDefault,
|
||||
// intrinsic?
|
||||
bool use_rdtscp_;
|
||||
|
||||
// The argument which corresponds to the profile counter buffer.
|
||||
llvm::Value* prof_counters_;
|
||||
|
||||
// The first read cycle counter in the program.
|
||||
llvm::Value* first_read_cycle_start_ = nullptr;
|
||||
|
||||
|
@ -593,7 +593,7 @@ class HloEvaluatorTypedVisitor : public DfsHloVisitorWithDefault {
|
||||
return Status::OK();
|
||||
}
|
||||
|
||||
Status HandleDivide(HloInstruction* divide) {
|
||||
Status HandleDivide(HloInstruction* divide) override {
|
||||
return HandleDivide<ElementwiseT>(divide);
|
||||
}
|
||||
|
||||
|
@ -2628,36 +2628,6 @@ Status HloInstruction::AcceptWithOperandOrder(
|
||||
return Status::OK();
|
||||
}
|
||||
|
||||
namespace {
|
||||
|
||||
// Returns true if the given order is a topological sort of the instructions
|
||||
// it contains.
|
||||
bool OrderIsTopologicalSort(const std::vector<const HloInstruction*>& order) {
|
||||
// Create a map from instruction to its position in 'order'.
|
||||
std::unordered_map<const HloInstruction*, int> order_position;
|
||||
for (int i = 0; i < order.size(); i++) {
|
||||
if (!order_position.insert({order[i], i}).second) {
|
||||
// Instruction order[i] is duplicated in the order.
|
||||
return false;
|
||||
}
|
||||
}
|
||||
// Verify that the operand of each instruction in the order is also in the
|
||||
// order *and* the operand's position is earlier (defs are before uses for
|
||||
// all ops).
|
||||
for (auto* instruction : order) {
|
||||
for (auto* operand : instruction->operands()) {
|
||||
if (!ContainsKey(order_position, operand) ||
|
||||
order_position.at(operand) >= order_position.at(instruction)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
Status HloInstruction::Accept(
|
||||
const std::function<Status(HloInstruction*)>& visitor_func) {
|
||||
FunctionVisitor visitor(visitor_func);
|
||||
|
@ -166,9 +166,6 @@ class HloValue : public BufferValue {
|
||||
|
||||
// Whether this value is live out of the HLO module.
|
||||
bool live_out_of_module_ = false;
|
||||
|
||||
// Whether this value is live out of its computation.
|
||||
bool live_out_of_computation_ = false;
|
||||
};
|
||||
|
||||
std::ostream& operator<<(std::ostream& out, const HloValue& hlo_value);
|
||||
|
@ -957,21 +957,6 @@ Status Service::TransferToClient(const TransferToClientRequest* arg,
|
||||
return Status::OK();
|
||||
}
|
||||
|
||||
namespace {
|
||||
|
||||
// Creates a clone of the given shaped buffer with the given device ordinal. The
|
||||
// shape and DeviceMemoryBase values of the clone are identical to the original.
|
||||
std::unique_ptr<ShapedBuffer> CloneShapedBufferOnDevice(
|
||||
const ShapedBuffer& shaped_buffer, int device_ordinal) {
|
||||
auto clone = absl::make_unique<ShapedBuffer>(
|
||||
shaped_buffer.on_host_shape(), shaped_buffer.on_device_shape(),
|
||||
shaped_buffer.platform(), device_ordinal);
|
||||
clone->buffers() = shaped_buffer.buffers();
|
||||
return clone;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
Status Service::TransferToServer(const TransferToServerRequest* arg,
|
||||
TransferToServerResponse* result) {
|
||||
TF_ASSIGN_OR_RETURN(Literal literal,
|
||||
|
@ -667,12 +667,11 @@ void ShapeTree<T>::CopySubtreeFrom(const ShapeTree<T>& other,
|
||||
template <typename T>
|
||||
bool ShapeTree<T>::operator==(const ShapeTree<T>& other) const {
|
||||
bool equal = true;
|
||||
ForEachElement(
|
||||
[this, &other, &equal](const ShapeIndex& index, const T& data) {
|
||||
if (data != other.element(index)) {
|
||||
equal = false;
|
||||
}
|
||||
});
|
||||
ForEachElement([&other, &equal](const ShapeIndex& index, const T& data) {
|
||||
if (data != other.element(index)) {
|
||||
equal = false;
|
||||
}
|
||||
});
|
||||
return equal;
|
||||
}
|
||||
|
||||
|
@ -37,6 +37,7 @@ limitations under the License.
|
||||
#include "tensorflow/core/platform/cpu_info.h"
|
||||
#include "tensorflow/core/platform/env.h"
|
||||
#include "tensorflow/core/platform/macros.h"
|
||||
#include "tensorflow/core/platform/mutex.h"
|
||||
#include "tensorflow/core/platform/types.h"
|
||||
|
||||
namespace xla {
|
||||
@ -756,10 +757,18 @@ class ShapeUtil {
|
||||
pool.emplace(tensorflow::Env::Default(), "foreach", kNumThreads);
|
||||
}
|
||||
|
||||
tensorflow::mutex mu;
|
||||
Status status; // Guarded by mu
|
||||
|
||||
while (n < rank) {
|
||||
if (pool != absl::nullopt) {
|
||||
pool->Schedule(
|
||||
[indexes, &visitor_function] { visitor_function(indexes); });
|
||||
pool->Schedule([indexes, &visitor_function, &mu, &status] {
|
||||
StatusOr<bool> result = visitor_function(indexes);
|
||||
if (!result.ok()) {
|
||||
tensorflow::mutex_lock lock(mu);
|
||||
status = status.ok() ? result.status() : status;
|
||||
}
|
||||
});
|
||||
} else {
|
||||
TF_ASSIGN_OR_RETURN(bool should_continue, visitor_function(indexes));
|
||||
if (!should_continue) {
|
||||
@ -777,7 +786,9 @@ class ShapeUtil {
|
||||
}
|
||||
}
|
||||
|
||||
return Status::OK();
|
||||
// Waits for the scheduled work to complete.
|
||||
pool.reset();
|
||||
return status;
|
||||
}
|
||||
|
||||
TF_DISALLOW_COPY_AND_ASSIGN(ShapeUtil);
|
||||
|
@ -511,7 +511,7 @@ void CollectiveParamResolverLocal::FindInstanceRec(
|
||||
if (irec->is_init) {
|
||||
exit_outside_locks = true;
|
||||
} else {
|
||||
irec->init_waiters.push_back([this, gr, cp, done](InstanceRec* irec) {
|
||||
irec->init_waiters.push_back([this, done](InstanceRec* irec) {
|
||||
CallbackWithStatus(done, irec);
|
||||
});
|
||||
return;
|
||||
|
@ -38,7 +38,7 @@ void CollectiveRemoteAccessLocal::RecvFromPeer(
|
||||
return;
|
||||
}
|
||||
buf_rendezvous_.ConsumeBuf(
|
||||
key, [this, to_tensor, to_device_ctx, to_device, to_alloc_attr,
|
||||
key, [to_tensor, to_device_ctx, to_device, to_alloc_attr,
|
||||
dev_to_dev_stream_index,
|
||||
done](const Status& s, BufRendezvous::Hook* hook) {
|
||||
if (!s.ok()) {
|
||||
|
@ -59,7 +59,7 @@ class EdgeSetTest;
|
||||
class Graph;
|
||||
class GraphDef;
|
||||
class Node;
|
||||
class OutputTensor;
|
||||
struct OutputTensor;
|
||||
class VersionDef;
|
||||
class WhileContext;
|
||||
|
||||
|
Loading…
Reference in New Issue
Block a user