message: [XLA][AMDGPU] use AddrSpaceCast to access GlobalVariable on AMDGPU
  In XLA kernels pointers are in default address space (0). On LLVM AMDGPU
  backend, all GlobalVariable instances must be in global address space (1).
  Therefore in an AddrSpaceCast is required for AMDGPU backend. To cope with
  existing NVPTX backend, use ConstantExpr::getPointerBitCastOrAddrSpaceCast
  in this commit.
diff: @@ -145,7 +145,7 @@ llvm::Value* HloToIrBindings::GetTypedIrValue(const HloInstruction& hlo,
  llvm::Value* typed_ir_value;
  if (llvm::isa<llvm::GlobalVariable>(ir_value)) {
  - typed_ir_value = llvm::ConstantExpr::getBitCast(
  + typed_ir_value = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(
  llvm::cast<llvm::GlobalVariable>(ir_value), dest_type);
  } else {
  typed_ir_value =
  