llvm/test/CodeGen/NVPTX/clmad.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mcpu=sm_100 -mattr=+ptx93 | FileCheck %s
; RUN: %if ptxas-sm_100 && ptxas-isa-9.3 %{ llc < %s -mcpu=sm_100 -mattr=+ptx93 | %ptxas-verify -arch=sm_100 %}
target triple = "nvptx64-nvidia-cuda"

define i64 @clmad_lo(i64 %a, i64 %b, i64 %c) {
; CHECK-LABEL: clmad_lo(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param::func.b64 %rd1, [clmad_lo_param_0];
; CHECK-NEXT:    ld.param::func.b64 %rd2, [clmad_lo_param_1];
; CHECK-NEXT:    ld.param::func.b64 %rd3, [clmad_lo_param_2];
; CHECK-NEXT:    clmad.lo.u64 %rd4, %rd1, %rd2, %rd3;
; CHECK-NEXT:    st.param::func.b64 [func_retval0], %rd4;
; CHECK-NEXT:    ret;
  %product = call i64 @llvm.clmul.i64(i64 %a, i64 %b)
  %result = xor i64 %product, %c
  ret i64 %result
}

define i64 @clmad_hi(i64 %a, i64 %b, i64 %c) {
; CHECK-LABEL: clmad_hi(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param::func.b64 %rd1, [clmad_hi_param_0];
; CHECK-NEXT:    ld.param::func.b64 %rd2, [clmad_hi_param_1];
; CHECK-NEXT:    ld.param::func.b64 %rd3, [clmad_hi_param_2];
; CHECK-NEXT:    clmad.hi.u64 %rd4, %rd1, %rd2, %rd3;
; CHECK-NEXT:    st.param::func.b64 [func_retval0], %rd4;
; CHECK-NEXT:    ret;
  %a.wide = zext i64 %a to i128
  %b.wide = zext i64 %b to i128
  %product = call i128 @llvm.clmul.i128(i128 %a.wide, i128 %b.wide)
  %shifted = lshr i128 %product, 64
  %high = trunc i128 %shifted to i64
  %result = xor i64 %high, %c
  ret i64 %result
}

define i64 @clmul_lo(i64 %a, i64 %b) {
; CHECK-LABEL: clmul_lo(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param::func.b64 %rd1, [clmul_lo_param_0];
; CHECK-NEXT:    ld.param::func.b64 %rd2, [clmul_lo_param_1];
; CHECK-NEXT:    clmad.lo.u64 %rd3, %rd1, %rd2, 0;
; CHECK-NEXT:    st.param::func.b64 [func_retval0], %rd3;
; CHECK-NEXT:    ret;
  %product = call i64 @llvm.clmul.i64(i64 %a, i64 %b)
  ret i64 %product
}

define i64 @clmul_hi(i64 %a, i64 %b) {
; CHECK-LABEL: clmul_hi(
; CHECK:       {
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param::func.b64 %rd1, [clmul_hi_param_0];
; CHECK-NEXT:    ld.param::func.b64 %rd2, [clmul_hi_param_1];
; CHECK-NEXT:    clmad.hi.u64 %rd3, %rd1, %rd2, 0;
; CHECK-NEXT:    st.param::func.b64 [func_retval0], %rd3;
; CHECK-NEXT:    ret;
  %a.wide = zext i64 %a to i128
  %b.wide = zext i64 %b to i128
  %product = call i128 @llvm.clmul.i128(i128 %a.wide, i128 %b.wide)
  %shifted = lshr i128 %product, 64
  %high = trunc i128 %shifted to i64
  ret i64 %high
}

declare i64 @llvm.clmul.i64(i64, i64)
declare i128 @llvm.clmul.i128(i128, i128)