// 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

    }

    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

less4:
	CMP   R5, $2
	BLT   less2
	MOVHZ (R4), R21
	SLD   R22, R21, R21
	OR    R16, R21, R16
	ADD   $16, R22
	ADD   $-2, R5
	ADD   $2, R4

less2:
	CMP   R5, $0
	BEQ   insert1
	MOVBZ (R4), R21
	SLD   R22, R21, R21
	OR    R16, R21, R16
	ADD   $8, R22

insert1:
	// Insert 1 at end of msg
	MOVD $1, R21
	SLD  R22, R21, R21
	OR   R16, R21, R16

carry: