#include "MurMurHash3.cuh"

__host__ __device__ inline uint64_t rotl64(uint64_t x, int8_t r)
{
  return (x << r) | (x >> (64 - r));
}

__host__ __device__ inline uint64_t getblock64(const uint64_t *p, int i)
{
  return p[i];
}

__host__ __device__ inline uint64_t fmix64(uint64_t k)
{
  k ^= k >> 33;
  k *= 0xff51afd7ed558ccd;
  k ^= k >> 33;
  k *= 0xc4ceb9fe1a85ec53;
  k ^= k >> 33;

  return k;
}

__host__ __device__ void MurmurHash3_128(const void *key, int len, uint32_t seed, void *out)
{
  const uint8_t *data = (const uint8_t *)key;
  const int nblocks = len / 16;

  uint64_t h1 = seed;
  uint64_t h2 = seed;

  const uint64_t c1 = 0x87c37b91114253d5;
  const uint64_t c2 = 0x4cf5ad432745937f;

  //----------
  // body

  const uint64_t *blocks = (const uint64_t *)(data);

  for (int i = 0; i < nblocks; i++)
  {
    uint64_t k1 = getblock64(blocks, i * 2 + 0);
    uint64_t k2 = getblock64(blocks, i * 2 + 1);

    k1 *= c1;
    k1 = rotl64(k1, 31);
    k1 *= c2;
    h1 ^= k1;

    h1 = rotl64(h1, 27);
    h1 += h2;
    h1 = h1 * 5 + 0x52dce729;

    k2 *= c2;
    k2 = rotl64(k2, 33);
    k2 *= c1;
    h2 ^= k2;

    h2 = rotl64(h2, 31);
    h2 += h1;
    h2 = h2 * 5 + 0x38495ab5;
  }

  //----------
  // tail

  const uint8_t *tail = (const uint8_t *)(data + nblocks * 16);

  uint64_t k1 = 0;
  uint64_t k2 = 0;

  switch (len & 15)
  {
  case 15:
    k2 ^= ((uint64_t)tail[14]) << 48;
  case 14:
    k2 ^= ((uint64_t)tail[13]) << 40;
  case 13:
    k2 ^= ((uint64_t)tail[12]) << 32;
  case 12:
    k2 ^= ((uint64_t)tail[11]) << 24;
  case 11:
    k2 ^= ((uint64_t)tail[10]) << 16;
  case 10:
    k2 ^= ((uint64_t)tail[9]) << 8;
  case 9:
    k2 ^= ((uint64_t)tail[8]) << 0;
    k2 *= c2;
    k2 = rotl64(k2, 33);
    k2 *= c1;
    h2 ^= k2;

  case 8:
    k1 ^= ((uint64_t)tail[7]) << 56;
  case 7:
    k1 ^= ((uint64_t)tail[6]) << 48;
  case 6:
    k1 ^= ((uint64_t)tail[5]) << 40;
  case 5:
    k1 ^= ((uint64_t)tail[4]) << 32;
  case 4:
    k1 ^= ((uint64_t)tail[3]) << 24;
  case 3:
    k1 ^= ((uint64_t)tail[2]) << 16;
  case 2:
    k1 ^= ((uint64_t)tail[1]) << 8;
  case 1:
    k1 ^= ((uint64_t)tail[0]) << 0;
    k1 *= c1;
    k1 = rotl64(k1, 31);
    k1 *= c2;
    h1 ^= k1;
  };

  //----------
  // finalization

  h1 ^= len;
  h2 ^= len;

  h1 += h2;
  h2 += h1;

  h1 = fmix64(h1);
  h2 = fmix64(h2);

  h1 += h2;
  h2 += h1;

  ((uint64_t *)out)[0] = h1;
  ((uint64_t *)out)[1] = h2;
}

	.file	"tmpxft_002db893_00000000-6_MurMurHash3.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2014:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2014:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text._Z6fmix64m,"axG",@progbits,_Z6fmix64m,comdat
	.weak	_Z6fmix64m
	.type	_Z6fmix64m, @function
_Z6fmix64m:
.LFB2010:
	.cfi_startproc
	endbr64
	movabsq	$-49064778989728563, %rdx
	movq	%rdi, %rax
	shrq	$33, %rax
	xorq	%rdi, %rax
	imulq	%rdx, %rax
	movq	%rax, %rdx
	shrq	$33, %rdx
	xorq	%rax, %rdx
	movabsq	$-4265267296055464877, %rax
	imulq	%rax, %rdx
	movq	%rdx, %rax
	shrq	$33, %rax
	xorq	%rdx, %rax
	ret
	.cfi_endproc
.LFE2010:
	.size	_Z6fmix64m, .-_Z6fmix64m
	.text
	.globl	_Z15MurmurHash3_128PKvijPv
	.type	_Z15MurmurHash3_128PKvijPv, @function
_Z15MurmurHash3_128PKvijPv:
.LFB2011:
	.cfi_startproc
	endbr64
	movl	%esi, %eax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r8
	xorl	%r9d, %r9d
	movabsq	$-8663945395140668459, %r10
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movl	%edx, %ebp
	cltd
	movabsq	$5545529020109919103, %r11
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rcx, %rbx
	movl	$16, %ecx
	idivl	%ecx
	movq	%rdi, %rcx
	movq	%rbp, %rdi
.L4:
	cmpl	%eax, %r9d
	jge	.L38
	movq	(%rcx), %rdx
	incl	%r9d
	addq	$16, %rcx
	imulq	%r10, %rdx
	rolq	$31, %rdx
	imulq	%r11, %rdx
	xorq	%rdi, %rdx
	rolq	$27, %rdx
	addq	%rbp, %rdx
	leaq	(%rdx,%rdx,4), %rdx
	leaq	1390208809(%rdx), %rdi
	movq	-8(%rcx), %rdx
	imulq	%r11, %rdx
	rorq	$31, %rdx
	imulq	%r10, %rdx
	xorq	%rbp, %rdx
	rolq	$31, %rdx
	addq	%rdi, %rdx
	leaq	(%rdx,%rdx,4), %rdx
	leaq	944331445(%rdx), %rbp
	jmp	.L4
.L38:
	sall	$4, %eax
	movslq	%eax, %rdx
	movl	%esi, %eax
	andl	$15, %eax
	addq	%r8, %rdx
	decl	%eax
	cmpl	$14, %eax
	ja	.L6
	leaq	.L8(%rip), %rcx
	movslq	(%rcx,%rax,4), %rax
	addq	%rcx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L8:
	.long	.L23-.L8
	.long	.L24-.L8
	.long	.L25-.L8
	.long	.L26-.L8
	.long	.L27-.L8
	.long	.L28-.L8
	.long	.L29-.L8
	.long	.L15-.L8
	.long	.L30-.L8
	.long	.L31-.L8
	.long	.L32-.L8
	.long	.L33-.L8
	.long	.L34-.L8
	.long	.L35-.L8
	.long	.L7-.L8
	.text
.L7:
	movzbl	14(%rdx), %eax
	salq	$48, %rax
	jmp	.L9
.L35:
	xorl	%eax, %eax
.L9:
	movzbl	13(%rdx), %ecx
	salq	$40, %rcx
	xorq	%rax, %rcx
	jmp	.L10
.L34:
	xorl	%ecx, %ecx
.L10:
	movzbl	12(%rdx), %eax
	salq	$32, %rax
	xorq	%rcx, %rax
	jmp	.L11
.L33:
	xorl	%eax, %eax
.L11:
	movzbl	11(%rdx), %ecx
	salq	$24, %rcx
	xorq	%rax, %rcx
	jmp	.L12
.L32:
	xorl	%ecx, %ecx
.L12:
	movzbl	10(%rdx), %eax
	salq	$16, %rax
	xorq	%rcx, %rax
	jmp	.L13
.L31:
	xorl	%eax, %eax
.L13:
	movzbl	9(%rdx), %ecx
	salq	$8, %rcx
	xorq	%rax, %rcx
	jmp	.L14
.L30:
	xorl	%ecx, %ecx
.L14:
	movzbl	8(%rdx), %eax
	xorq	%rcx, %rax
	movabsq	$5545529020109919103, %rcx
	imulq	%rcx, %rax
	movabsq	$-8663945395140668459, %rcx
	rorq	$31, %rax
	imulq	%rcx, %rax
	xorq	%rax, %rbp
.L15:
	movzbl	7(%rdx), %ecx
	salq	$56, %rcx
	jmp	.L16
.L29:
	xorl	%ecx, %ecx
.L16:
	movzbl	6(%rdx), %eax
	salq	$48, %rax
	xorq	%rcx, %rax
	jmp	.L17
.L28:
	xorl	%eax, %eax
.L17:
	movzbl	5(%rdx), %ecx
	salq	$40, %rcx
	xorq	%rax, %rcx
	jmp	.L18
.L27:
	xorl	%ecx, %ecx
.L18:
	movzbl	4(%rdx), %eax
	salq	$32, %rax
	xorq	%rcx, %rax
	jmp	.L19
.L26:
	xorl	%eax, %eax
.L19:
	movzbl	3(%rdx), %ecx
	salq	$24, %rcx
	xorq	%rax, %rcx
	jmp	.L20
.L25:
	xorl	%ecx, %ecx
.L20:
	movzbl	2(%rdx), %eax
	salq	$16, %rax
	xorq	%rcx, %rax
	jmp	.L21
.L24:
	xorl	%eax, %eax
.L21:
	movzbl	1(%rdx), %ecx
	salq	$8, %rcx
	xorq	%rax, %rcx
	jmp	.L22
.L23:
	xorl	%ecx, %ecx
.L22:
	movzbl	(%rdx), %eax
	movabsq	$-8663945395140668459, %rdx
	xorq	%rcx, %rax
	imulq	%rdx, %rax
	movabsq	$5545529020109919103, %rdx
	rolq	$31, %rax
	imulq	%rdx, %rax
	xorq	%rax, %rdi
.L6:
	movslq	%esi, %rsi
	xorq	%rsi, %rdi
	xorq	%rbp, %rsi
	addq	%rsi, %rdi
	leaq	(%rsi,%rdi), %rbp
	call	_Z6fmix64m
	movq	%rbp, %rdi
	movq	%rax, %r12
	call	_Z6fmix64m
	leaq	(%r12,%rax), %rdx
	addq	%rdx, %rax
	movq	%rdx, (%rbx)
	movq	%rax, 8(%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2011:
	.size	_Z15MurmurHash3_128PKvijPv, .-_Z15MurmurHash3_128PKvijPv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2037:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

/**
 * @file
 * @brief The MurMur3 128-bit hash, with CUDA support
 * @see https://github.com/aappleby/smhasher
 */
#ifndef MURMUR_HASH_3_CUH_
#define MURMUR_HASH_3_CUH_

#include <cstdint>

/**
 * Create a 128-bit hash
 *
 * Actually MurmurHash3_x64_128
 *
 * @example
 *   char *elem = "aaa123";
 *   uint64_t hash[2];
 *   MurmurHash3_128(elem, sizeof(elem), 0, &hash);
 *   // hash[0] and hash[1] contain the hash
 *
 * @param key The pointer to the value to be hashed
 * @param len The length of the key
 * @param seed Any 32-bit value
 * @param out The pointer to the output array
 */
__host__ __device__ void MurmurHash3_128(const void *key, int len, uint32_t seed, void *out);

#endif // MURMUR_HASH_3_CUH_

	.text
	.file	"MurMurHash3.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
