// license that can be found in the LICENSE file.

//go:build !purego

#include "textflag.h"

// func block(dig *digest, p []byte)
TEXT ·block(SB), NOSPLIT|NOFRAME, $0-32
	MOVBZ  ·useAsm(SB), R4
	LMG    dig+0(FP), R1, R3            // R2 = &p[0], R3 = len(p)
	MOVBZ  $2, R0                       // SHA-256 function code
	CMPBEQ R4, $0, generic

loop:
	KIMD R0, R2      // compute intermediate message digest (KIMD)

// license that can be found in the LICENSE file.

//go:build !purego

#include "textflag.h"

// func block(dig *digest, p []byte)
TEXT ·block(SB), NOSPLIT|NOFRAME, $0-32
	MOVBZ  ·useAsm(SB), R4
	LMG    dig+0(FP), R1, R3            // R2 = &p[0], R3 = len(p)
	MOVBZ  $3, R0                       // SHA-512 function code
	CMPBEQ R4, $0, generic

loop:
	KIMD R0, R2      // compute intermediate message digest (KIMD)

        "The device tensor has no associated device buffer."));
    return;
  }

  xla::PjRtFuture<> future = device_buffer->ToLiteral(literal.get());
  future.OnReady([literal = std::move(literal), done = std::move(done)](
                     const tensorflow::Status& status) { done(status); });
}

void PjRtDeviceContext::CopyCPUTensorToDevice(const Tensor* cpu_tensor,

      new stream_executor::CPlatform(
          std::move(platform), params.destroy_platform, std::move(platform_fns),
          params.destroy_platform_fns, std::move(device_fns), std::move(se),
          std::move(timer_fns)));
  TF_CHECK_OK(
      stream_executor::PlatformManager::RegisterPlatform(std::move(cplatform)));
  // TODO(annarev): Return `use_bfc_allocator` value in some way so that it is

class ReproducerStream : public mlir::ReproducerStream {
 public:
  ReproducerStream(std::string name, std::unique_ptr<llvm::raw_ostream> os)
      : name_(std::move(name)), os_(std::move(os)) {}

  llvm::StringRef description() override { return name_; }

  llvm::raw_ostream& os() override { return *os_; }

 private:
  std::string name_;
  std::unique_ptr<llvm::raw_ostream> os_;

    }

    llvm::SourceMgr source_mgr;
    source_mgr.AddNewSourceBuffer(std::move(file), llvm::SMLoc());
    auto module = parseSourceFile<ModuleOp>(source_mgr, context);
    if (module->getOperation() == nullptr) {
      return absl::AbortedError("Failed to parse input MLIR model.");
    }

    return quant::stablehlo::ImportedMlirModuleOp(std::move(module), nullptr);
  }

#define MSGSCHEDULE0(index) \
	MOVBU	((index*8)+0)(X29), X5; \
	MOVBU	((index*8)+1)(X29), X6; \
	MOVBU	((index*8)+2)(X29), X7; \
	MOVBU	((index*8)+3)(X29), X8; \
	SLL	$56, X5; \
	SLL	$48, X6; \
	OR	X5, X6, X5; \
	SLL	$40, X7; \
	OR	X5, X7, X5; \
	SLL	$32, X8; \
	OR	X5, X8, X5; \
	MOVBU	((index*8)+4)(X29), X9; \
	MOVBU	((index*8)+5)(X29), X6; \
	MOVBU	((index*8)+6)(X29), X7; \
	MOVBU	((index*8)+7)(X29), X8; \

//go:build linux && (ppc64 || ppc64le)

#include "textflag.h"

// func Syscall6(num, a1, a2, a3, a4, a5, a6 uintptr) (r1, r2, errno uintptr)
TEXT ·Syscall6<ABIInternal>(SB),NOSPLIT,$0-80
	MOVD	R3, R10	// Move syscall number to R10. SYSCALL will move it R0, and restore R0.
	MOVD	R4, R3
	MOVD	R5, R4
	MOVD	R6, R5
	MOVD	R7, R6
	MOVD	R8, R7
	MOVD	R9, R8
	SYSCALL	R10
	MOVD	$-1, R6
	ISEL	CR0SO, R3, R0, R5 // errno = (error) ? R3 : 0

    SessionOptions options;
    std::vector<std::unique_ptr<Device>> devices;
    Status s = device_factory->CreateDevices(
        options, "/job:worker/replica:0/task:0", &devices);
    device_ = std::move(devices[0]);

    AllocatorAttributes host_alloc_attr;
    host_alloc_attr.set_on_host(true);
    host_allocator_ = device_->GetAllocator(host_alloc_attr);

    AllocatorAttributes device_alloc_attr;

//go:linkname Malloc runtime.msanmalloc
func Malloc(addr unsafe.Pointer, sz uintptr)

//go:linkname Free runtime.msanfree
func Free(addr unsafe.Pointer, sz uintptr)

//go:linkname Move runtime.msanmove