clang/test/CodeGen/builtins-nvptx-native-half-type.c

   1 // REQUIRES: nvptx-registered-target
   2 //
   3 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
   4 // RUN:   sm_75 -target-feature +ptx70 -fcuda-is-device -fnative-half-type \
   5 // RUN:   -emit-llvm -o - -x cuda %s \
   6 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM75 %s
   7
   8 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
   9 // RUN:   sm_80 -target-feature +ptx70 -fcuda-is-device -fnative-half-type \
  10 // RUN:   -emit-llvm -o - -x cuda %s \
  11 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 %s
  12
  13 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \
  14 // RUN:   -target-cpu sm_80 -target-feature +ptx70 -fcuda-is-device \
  15 // RUN:   -fnative-half-type -emit-llvm -o - -x cuda %s \
  16 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 %s
  17
  18 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
  19 // RUN:   sm_86 -target-feature +ptx72 -fcuda-is-device -fnative-half-type \
  20 // RUN:   -emit-llvm -o - -x cuda %s \
  21 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 %s
  22
  23 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \
  24 // RUN:   -target-cpu sm_86 -target-feature +ptx72 -fcuda-is-device \
  25 // RUN:   -fnative-half-type -emit-llvm -o - -x cuda %s \
  26 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX72_SM86 %s
  27
  28 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
  29 // RUN:   sm_53 -target-feature +ptx42 -fcuda-is-device -fnative-half-type \
  30 // RUN:   -emit-llvm -o - -x cuda %s \
  31 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX42_SM53 %s
  32
  33 // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \
  34 // RUN:   -target-cpu sm_53 -target-feature +ptx42 -fcuda-is-device \
  35 // RUN:   -fnative-half-type -emit-llvm -o - -x cuda %s \
  36 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX42_SM53 %s
  37
  38 #define __device__ __attribute__((device))
  39
  40 __device__ void nvvm_ex2_sm75() {
  41 #if __CUDA_ARCH__ >= 750
  42   // CHECK_PTX70_SM75: call half @llvm.nvvm.ex2.approx.f16
  43   __nvvm_ex2_approx_f16(0.1f16);
  44   // CHECK_PTX70_SM75: call <2 x half> @llvm.nvvm.ex2.approx.f16x2
  45   __nvvm_ex2_approx_f16x2({0.1f16, 0.7f16});
  46 #endif
  47   // CHECK: ret void
  48 }
  49
  50 // CHECK-LABEL: nvvm_min_max_sm80
  51 __device__ void nvvm_min_max_sm80() {
  52 #if __CUDA_ARCH__ >= 800
  53   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.f16
  54   __nvvm_fmin_f16(0.1f16, 0.1f16);
  55   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.ftz.f16
  56   __nvvm_fmin_ftz_f16(0.1f16, 0.1f16);
  57   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.nan.f16
  58   __nvvm_fmin_nan_f16(0.1f16, 0.1f16);
  59   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.ftz.nan.f16
  60   __nvvm_fmin_ftz_nan_f16(0.1f16, 0.1f16);
  61   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.f16x2
  62   __nvvm_fmin_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  63   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.ftz.f16x2
  64   __nvvm_fmin_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  65   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.nan.f16x2
  66   __nvvm_fmin_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  67   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.ftz.nan.f16x2
  68   __nvvm_fmin_ftz_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  69
  70   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.f16
  71   __nvvm_fmax_f16(0.1f16, 0.1f16);
  72   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.ftz.f16
  73   __nvvm_fmax_ftz_f16(0.1f16, 0.1f16);
  74   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.nan.f16
  75   __nvvm_fmax_nan_f16(0.1f16, 0.1f16);
  76   // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.ftz.nan.f16
  77   __nvvm_fmax_ftz_nan_f16(0.1f16, 0.1f16);
  78   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.f16x2
  79   __nvvm_fmax_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  80   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.ftz.f16x2
  81   __nvvm_fmax_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  82   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.nan.f16x2
  83   __nvvm_fmax_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  84   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.ftz.nan.f16x2
  85   __nvvm_fmax_ftz_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
  86 #endif
  87   // CHECK: ret void
  88 }
  89
  90 // CHECK-LABEL: nvvm_fma_f16_f16x2_sm80
  91 __device__ void nvvm_fma_f16_f16x2_sm80() {
  92 #if __CUDA_ARCH__ >= 800
  93   // CHECK_PTX70_SM80: call half @llvm.nvvm.fma.rn.relu.f16
  94   __nvvm_fma_rn_relu_f16(0.1f16, 0.1f16, 0.1f16);
  95   // CHECK_PTX70_SM80: call half @llvm.nvvm.fma.rn.ftz.relu.f16
  96   __nvvm_fma_rn_ftz_relu_f16(0.1f16, 0.1f16, 0.1f16);
  97
  98   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fma.rn.relu.f16x2
  99   __nvvm_fma_rn_relu_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
 100                            {0.1f16, 0.7f16});
 101   // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fma.rn.ftz.relu.f16x2
 102   __nvvm_fma_rn_ftz_relu_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
 103                                {0.1f16, 0.7f16});
 104 #endif
 105   // CHECK: ret void
 106 }
 107
 108 // CHECK-LABEL: nvvm_fma_f16_f16x2_sm53
 109 __device__ void nvvm_fma_f16_f16x2_sm53() {
 110 #if __CUDA_ARCH__ >= 530
 111   // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.f16
 112   __nvvm_fma_rn_f16(0.1f16, 0.1f16, 0.1f16);
 113   // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.ftz.f16
 114   __nvvm_fma_rn_ftz_f16(0.1f16, 0.1f16, 0.1f16);
 115   // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.sat.f16
 116   __nvvm_fma_rn_sat_f16(0.1f16, 0.1f16, 0.1f16);
 117   // CHECK_PTX42_SM53: call half @llvm.nvvm.fma.rn.ftz.sat.f16
 118   __nvvm_fma_rn_ftz_sat_f16(0.1f16, 0.1f16, 0.1f16);
 119
 120   // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.f16x2
 121   __nvvm_fma_rn_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
 122                       {0.1f16, 0.7f16});
 123   // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2
 124   __nvvm_fma_rn_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
 125                           {0.1f16, 0.7f16});
 126   // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2
 127   __nvvm_fma_rn_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
 128                           {0.1f16, 0.7f16});
 129   // CHECK_PTX42_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2
 130   __nvvm_fma_rn_ftz_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},
 131                               {0.1f16, 0.7f16});
 132 #endif
 133   // CHECK: ret void
 134 }
 135
 136 // CHECK-LABEL: nvvm_min_max_sm86
 137 __device__ void nvvm_min_max_sm86() {
 138 #if __CUDA_ARCH__ >= 860
 139   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.xorsign.abs.f16
 140   __nvvm_fmin_xorsign_abs_f16(0.1f16, 0.1f16);
 141   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.ftz.xorsign.abs.f16
 142   __nvvm_fmin_ftz_xorsign_abs_f16(0.1f16, 0.1f16);
 143   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.nan.xorsign.abs.f16
 144   __nvvm_fmin_nan_xorsign_abs_f16(0.1f16, 0.1f16);
 145   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16
 146   __nvvm_fmin_ftz_nan_xorsign_abs_f16(0.1f16, 0.1f16);
 147   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmin.xorsign.abs.f16x2
 148   __nvvm_fmin_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 149   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmin.ftz.xorsign.abs.f16x2
 150   __nvvm_fmin_ftz_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 151   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmin.nan.xorsign.abs.f16x2
 152   __nvvm_fmin_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 153   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16x2
 154   __nvvm_fmin_ftz_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 155
 156   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.xorsign.abs.f16
 157   __nvvm_fmax_xorsign_abs_f16(0.1f16, 0.1f16);
 158   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.ftz.xorsign.abs.f16
 159   __nvvm_fmax_ftz_xorsign_abs_f16(0.1f16, 0.1f16);
 160   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.nan.xorsign.abs.f16
 161   __nvvm_fmax_nan_xorsign_abs_f16(0.1f16, 0.1f16);
 162   // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16
 163   __nvvm_fmax_ftz_nan_xorsign_abs_f16(0.1f16, 0.1f16);
 164   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.xorsign.abs.f16x2
 165   __nvvm_fmax_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 166   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.ftz.xorsign.abs.f16x2
 167   __nvvm_fmax_ftz_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 168   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.nan.xorsign.abs.f16x2
 169   __nvvm_fmax_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 170   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16x2
 171   __nvvm_fmax_ftz_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 172 #endif
 173   // CHECK: ret void
 174 }
 175
 176 typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2)));
 177
 178 // CHECK-LABEL: nvvm_ldg_native_half_types
 179 __device__ void nvvm_ldg_native_half_types(const void *p) {
 180   // CHECK: load half, ptr addrspace(1) {{.*}}, align 2, !invariant.load
 181   __nvvm_ldg_h((const __fp16 *)p);
 182   // CHECK: load <2 x half>, ptr addrspace(1) {{.*}}, align 4, !invariant.load
 183   __nvvm_ldg_h2((const __fp16v2 *)p);
 184 }
 185
 186 // CHECK-LABEL: nvvm_ldu_native_half_types
 187 __device__ void nvvm_ldu_native_half_types(const void *p) {
 188   // CHECK: call half @llvm.nvvm.ldu.global.f.f16.p0
 189   __nvvm_ldu_h((const __fp16 *)p);
 190   // CHECK: call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p0
 191   __nvvm_ldu_h2((const __fp16v2 *)p);
 192 }