[XLA] Rename OPAQUE to OPAQUE_TYPE.

The identifier `OPAQUE` is used as a macro in windows.h.

This should (at least partially) fix #25773.

Textual HLO remains unchanged (i.e. it still uses "opaque").  The parser now
accepts either "opaque_type" or "opaque".

PiperOrigin-RevId: 245055351

tensorflow/compiler/tf2xla/lib/util.cc

@@ -111,7 +111,7 @@ xla::XlaOp IntegerLiteral(xla::XlaBuilder* builder, xla::PrimitiveType type,
       break;
     case xla::TUPLE:
       LOG(FATAL) << "tuple element type is not integral";
-    case xla::OPAQUE:
+    case xla::OPAQUE_TYPE:
       LOG(FATAL) << "opaque element type is not integral";
     default:
       LOG(FATAL) << "unhandled element type " << type;

tensorflow/compiler/xla/client/lib/math_test.cc

@@ -180,7 +180,7 @@ XLA_TEST_F(MathTest, RealFpOnlyOps) {
       shape = ShapeUtil::MakeShape(ty, {42});
     } else if (ty == PrimitiveType::TUPLE) {
       shape = ShapeUtil::MakeTupleShape({});
-    } else if (ty == PrimitiveType::OPAQUE) {
+    } else if (ty == PrimitiveType::OPAQUE_TYPE) {
       shape = ShapeUtil::MakeOpaqueShape();
     } else if (ty == PrimitiveType::TOKEN) {
       shape = ShapeUtil::MakeTokenShape();

tensorflow/compiler/xla/literal.cc

@@ -293,8 +293,9 @@ Status MutableLiteralBase::CopyElementFrom(const LiteralSlice& src_literal,
     return InvalidArgument("LiteralProto has no shape");
   }
   Shape shape(proto.shape());
-  if (ShapeUtil::HasPrimitiveType(shape, OPAQUE)) {
-    return InvalidArgument("Literal shape cannot include OPAQUE sub-shape");
+  if (ShapeUtil::HasPrimitiveType(shape, OPAQUE_TYPE)) {
+    return InvalidArgument(
+        "Literal shape cannot include OPAQUE_TYPE sub-shape");
   }
   if (!LayoutUtil::HasLayout(shape)) {
     return InvalidArgument("LiteralProto has no layout");

tensorflow/compiler/xla/literal_util.cc

@@ -136,7 +136,7 @@ Literal ConvertType(LiteralSlice literal) {
       return LiteralUtil::CreateR0<bool>(false);
     case TUPLE:
       LOG(FATAL) << "tuple element type cannot take on value of 0";
-    case OPAQUE:
+    case OPAQUE_TYPE:
       LOG(FATAL) << "opaque element type cannot take on value of 0";
     default:
       LOG(FATAL) << "Unhandled primitive type " << primitive_type;
@@ -176,7 +176,7 @@ Literal ConvertType(LiteralSlice literal) {
       LOG(FATAL) << "u16/s16 literals not yet implemented";
     case TUPLE:
       LOG(FATAL) << "tuple element type cannot take on value of 1";
-    case OPAQUE:
+    case OPAQUE_TYPE:
       LOG(FATAL) << "opaque element type cannot take on value of 1";
     default:
       LOG(FATAL) << "Unhandled primitive type " << primitive_type;
@@ -220,7 +220,7 @@ Literal ConvertType(LiteralSlice literal) {
           static_cast<bfloat16>(-std::numeric_limits<float>::infinity()));
     case TUPLE:
       LOG(FATAL) << "tuple element type has no minimum value";
-    case OPAQUE:
+    case OPAQUE_TYPE:
       LOG(FATAL) << "opaque element type has no minimum value";
     default:
       LOG(FATAL) << "Unhandled primitive type " << primitive_type;
@@ -260,7 +260,7 @@ Literal ConvertType(LiteralSlice literal) {
           static_cast<bfloat16>(std::numeric_limits<float>::infinity()));
     case TUPLE:
       LOG(FATAL) << "tuple element type has no maximum value";
-    case OPAQUE:
+    case OPAQUE_TYPE:
       LOG(FATAL) << "opaque element type has no maximum value";
     default:
       LOG(FATAL) << "Unhandled primitive type " << primitive_type;

tensorflow/compiler/xla/primitive_util.cc

@@ -89,8 +89,8 @@ int BitWidth(PrimitiveType type) {
     case TUPLE:
       LOG(FATAL) << "TUPLE is an invalid type for BitWidth";
 
-    case OPAQUE:
-      LOG(FATAL) << "OPAQUE is an invalid type for BitWidth";
+    case OPAQUE_TYPE:
+      LOG(FATAL) << "OPAQUE_TYPE is an invalid type for BitWidth";
 
     default:
       LOG(FATAL) << "Unhandled primitive type " << type;
@@ -126,17 +126,22 @@ PrimitiveType ComplexComponentType(PrimitiveType complex_type) {
 
 bool IsArrayType(PrimitiveType primitive_type) {
   return primitive_type != PRIMITIVE_TYPE_INVALID && primitive_type != TUPLE &&
-         primitive_type != OPAQUE && primitive_type != TOKEN;
+         primitive_type != OPAQUE_TYPE && primitive_type != TOKEN;
 }
 
 // Class to memoize the computation of
 //   absl::AsciiStrToLower(PrimitiveType_Name(p))
 // for all PrimitiveType values "p"
+//
+// xla::OPAQUE_TYPE canonically maps to the string "opaque" -- the only reason
+// it's called OPAQUE_TYPE is to avoid clashing with a windows.h macro.
 class PrimitiveTypeNameGenerator {
  public:
   PrimitiveTypeNameGenerator() {
     for (int i = 0; i < PrimitiveType_ARRAYSIZE; i++) {
-      if (PrimitiveType_IsValid(i)) {
+      if (i == static_cast<int>(OPAQUE_TYPE)) {
+        lowercase_name_[i] = "opaque";
+      } else if (PrimitiveType_IsValid(i)) {
         lowercase_name_[i] = absl::AsciiStrToLower(
             PrimitiveType_Name(static_cast<PrimitiveType>(i)));
       }
@@ -158,6 +163,9 @@ const string& LowercasePrimitiveTypeName(PrimitiveType s) {
 namespace {
 
 // Returns a map from lower-case primitive type name to primitive type.
+//
+// Due to Postel's Law considerations, both "opaque" and "opaque_type" map to
+// the xla::OPAQUE_TYPE enumerator.
 const std::unordered_map<string, PrimitiveType>& GetPrimitiveTypeStringMap() {
   static std::unordered_map<string, PrimitiveType>* name_to_type = [] {
     static auto* map = new std::unordered_map<string, PrimitiveType>;
@@ -167,6 +175,7 @@ const std::unordered_map<string, PrimitiveType>& GetPrimitiveTypeStringMap() {
         (*map)[LowercasePrimitiveTypeName(value)] = value;
       }
     }
+    (*map)["opaque"] = OPAQUE_TYPE;
     return map;
   }();
   return *name_to_type;

tensorflow/compiler/xla/python/xla.cc

@@ -122,7 +122,7 @@ PYBIND11_MODULE(xla_extension, m) {
       .value("C64", C64)
       .value("C128", C128)
       .value("TUPLE", TUPLE)
-      .value("OPAQUE", OPAQUE)
+      .value("OPAQUE_TYPE", OPAQUE_TYPE)
       .value("TOKEN", TOKEN);
 
   // Shapes

tensorflow/compiler/xla/service/hlo_evaluator.cc

@@ -215,10 +215,10 @@ HloEvaluator::HloEvaluator(int64 max_loop_iterations)
         return Unimplemented(
             "HloEvaluatorTypedVisitor: unhandled primitive type: TUPLE.");
       });
-  typed_visitors_[OPAQUE] =
+  typed_visitors_[OPAQUE_TYPE] =
       absl::make_unique<FunctionVisitor>([](HloInstruction*) {
         return Unimplemented(
-            "HloEvaluatorTypedVisitor: unhandled primitive type: OPAQUE.");
+            "HloEvaluatorTypedVisitor: unhandled primitive type: OPAQUE_TYPE.");
       });
   typed_visitors_[TOKEN] =
       absl::make_unique<FunctionVisitor>([](HloInstruction*) {
@@ -497,7 +497,7 @@ Status HloEvaluator::HandleIsFinite(HloInstruction* is_finite) {
   switch (elem_ty) {
     case PRED:
     case TUPLE:
-    case OPAQUE:
+    case OPAQUE_TYPE:
     case TOKEN:
     case S8:
     case S16:

tensorflow/compiler/xla/service/hlo_graph_dumper.cc

@@ -117,7 +117,7 @@ class NodeFilter {
 // We arbitrarily set this as the boundary between "large" and "small"
 // instructions.
 bool IsSmall(const HloInstruction* instr) {
-  if (ShapeUtil::HasPrimitiveType(instr->shape(), OPAQUE) ||
+  if (ShapeUtil::HasPrimitiveType(instr->shape(), OPAQUE_TYPE) ||
       ShapeUtil::HasPrimitiveType(instr->shape(), TOKEN)) {
     return true;
   }

tensorflow/compiler/xla/service/llvm_ir/llvm_util.cc

@@ -194,7 +194,7 @@ llvm::Type* PrimitiveTypeToIrType(PrimitiveType element_type,
     }  // A Tuple contains an array of pointers. Use i8*.
     case TUPLE:
     // An Opaque is like a void*, use i8*.
-    case OPAQUE:
+    case OPAQUE_TYPE:
       return llvm::Type::getInt8PtrTy(module->getContext());
     case TOKEN:
       // Tokens do not have a physical representation, but the compiler needs

tensorflow/compiler/xla/shape.h

@@ -56,7 +56,7 @@ class Shape {
   bool IsArray() const { return primitive_util::IsArrayType(element_type()); }
   bool IsTuple() const { return element_type() == TUPLE; }
   bool IsToken() const { return element_type() == TOKEN; }
-  bool IsOpaque() const { return element_type() == OPAQUE; }
+  bool IsOpaque() const { return element_type() == OPAQUE_TYPE; }
 
   // Returns true if no array dimension in the shape is dynamically sized. Tuple
   // shapes are traversed recursively.

tensorflow/compiler/xla/shape_util.cc

@@ -96,7 +96,7 @@ StatusOr<Shape> MakeShapeWithLayoutInternal(
     return InvalidArgument("Dimensions size is %ld, but layout size is %ld.",
                            dimensions.size(), minor_to_major.size());
   }
-  if (element_type == OPAQUE || element_type == TUPLE) {
+  if (element_type == OPAQUE_TYPE || element_type == TUPLE) {
     return InvalidArgument("Unsupported element type: %s",
                            PrimitiveType_Name(element_type));
   }
@@ -258,7 +258,7 @@ ShapeUtil::MakeShapeWithDescendingLayoutAndSamePhysicalLayout(
 
 /* static */ Shape ShapeUtil::MakeOpaqueShape() {
   Shape result;
-  result.set_element_type(OPAQUE);
+  result.set_element_type(OPAQUE_TYPE);
   TF_DCHECK_OK(ValidateShapeWithOptionalLayout(result));
   return result;
 }
@@ -319,7 +319,7 @@ ShapeUtil::MakeShapeWithDescendingLayoutAndSamePhysicalLayout(
     case C64:
     case C128:
     case TUPLE:
-    case OPAQUE:
+    case OPAQUE_TYPE:
     case TOKEN:
       return false;
 
@@ -570,7 +570,7 @@ ShapeUtil::MakeShapeWithDescendingLayoutAndSamePhysicalLayout(
       // Tokens require no space.
       return 0;
     case TUPLE:
-    case OPAQUE:
+    case OPAQUE_TYPE:
       LOG(FATAL) << PrimitiveType_Name(primitive_type)
                  << " primitive type has no definitive size";
     default:
@@ -591,7 +591,7 @@ ShapeUtil::MakeShapeWithDescendingLayoutAndSamePhysicalLayout(
     return byte_size;
   } else if (shape.element_type() == TOKEN) {
     return 0;
-  } else if (shape.element_type() == OPAQUE) {
+  } else if (shape.element_type() == OPAQUE_TYPE) {
     CHECK_GT(pointer_size, 0);
     return pointer_size;
   }
@@ -653,7 +653,7 @@ ShapeUtil::MakeShapeWithDescendingLayoutAndSamePhysicalLayout(
   }
 
   // Tokens and opaques can should not have layout or dimensions.
-  if (shape.element_type() == TOKEN || shape.element_type() == OPAQUE) {
+  if (shape.element_type() == TOKEN || shape.element_type() == OPAQUE_TYPE) {
     if (shape.dimensions_size() != 0) {
       return InvalidArgument(
           "shape has %s element type, but has dimensions field: %s",

tensorflow/compiler/xla/xla_data.proto

@@ -71,7 +71,10 @@ enum PrimitiveType {
   // An opaque type used for passing context-specific data to a custom
   // operation. Shapes of this primitive type will have empty dimensions and
   // tuple_shapes fields.
-  OPAQUE = 14;
+  //
+  // (OPAQUE would be a better name for this identifier, but that conflicts with
+  // a macro defined in windows.h.)
+  OPAQUE_TYPE = 14;
 
   // A token type threaded between side-effecting operations. Shapes of this
   // primitive type will have empty dimensions and tuple_shapes fields.