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

   1 // REQUIRES: nvptx-registered-target
   2 //
   3 // RUN: not %clang_cc1 -fsyntax-only -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
   4 // RUN:   sm_86 -target-feature +ptx72 -fcuda-is-device -x cuda -emit-llvm -o - %s 2>&1 \
   5 // RUN:   | FileCheck -check-prefix=CHECK_ERROR %s
   6
   7 #define __device__ __attribute__((device))
   8 typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2)));
   9
  10 __device__ void nvvm_native_half_types(void *a, void*b, void*c, __fp16* out) {
  11   __fp16v2 resv2 = {0, 0};
  12   *out += __nvvm_ex2_approx_f16(*(__fp16 *)a);
  13   resv2 = __nvvm_ex2_approx_f16x2(*(__fp16v2*)a);
  14
  15   *out += __nvvm_fma_rn_relu_f16(*(__fp16*)a, *(__fp16*)b, *(__fp16*)c);
  16   *out += __nvvm_fma_rn_ftz_relu_f16(*(__fp16*)a, *(__fp16*)b, *(__fp16 *)c);
  17   resv2 += __nvvm_fma_rn_relu_f16x2(*(__fp16v2*)a, *(__fp16v2*)b, *(__fp16v2*)c);
  18   resv2 += __nvvm_fma_rn_ftz_relu_f16x2(*(__fp16v2*)a, *(__fp16v2*)b, *(__fp16v2*)c);
  19   *out += __nvvm_fma_rn_ftz_f16(*(__fp16*)a, *(__fp16*)b, *(__fp16*)c);
  20   *out += __nvvm_fma_rn_sat_f16(*(__fp16*)a, *(__fp16*)b, *(__fp16*)c);
  21   *out += __nvvm_fma_rn_ftz_sat_f16(*(__fp16*)a, *(__fp16*)b, *(__fp16*)c);
  22   resv2 += __nvvm_fma_rn_f16x2(*(__fp16v2*)a, *(__fp16v2*)b, *(__fp16v2*)c);
  23   resv2 += __nvvm_fma_rn_ftz_f16x2(*(__fp16v2*)a, *(__fp16v2*)b, *(__fp16v2*)c);
  24   resv2 += __nvvm_fma_rn_sat_f16x2(*(__fp16v2*)a, *(__fp16v2*)b, *(__fp16v2*)c);
  25   resv2 += __nvvm_fma_rn_ftz_sat_f16x2(*(__fp16v2*)a, *(__fp16v2*)b, *(__fp16v2*)c);
  26
  27   *out += __nvvm_fmin_f16(*(__fp16*)a, *(__fp16*)b);
  28   *out += __nvvm_fmin_ftz_f16(*(__fp16*)a, *(__fp16*)b);
  29   *out += __nvvm_fmin_nan_f16(*(__fp16*)a, *(__fp16*)b);
  30   *out += __nvvm_fmin_ftz_nan_f16(*(__fp16*)a, *(__fp16*)b);
  31   resv2 += __nvvm_fmin_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  32   resv2 += __nvvm_fmin_ftz_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  33   resv2 += __nvvm_fmin_nan_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  34   resv2 += __nvvm_fmin_ftz_nan_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  35   *out += __nvvm_fmin_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  36   *out += __nvvm_fmin_ftz_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  37   *out += __nvvm_fmin_nan_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  38   *out += __nvvm_fmin_ftz_nan_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  39   resv2 += __nvvm_fmin_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  40   resv2 += __nvvm_fmin_ftz_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  41   resv2 += __nvvm_fmin_nan_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  42   resv2 += __nvvm_fmin_ftz_nan_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  43
  44   *out += __nvvm_fmax_f16(*(__fp16*)a, *(__fp16*)b);
  45   *out += __nvvm_fmax_ftz_f16(*(__fp16*)a, *(__fp16*)b);
  46   *out += __nvvm_fmax_nan_f16(*(__fp16*)a, *(__fp16*)b);
  47   *out += __nvvm_fmax_ftz_nan_f16(*(__fp16*)a, *(__fp16*)b);
  48   resv2 += __nvvm_fmax_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  49   resv2 += __nvvm_fmax_ftz_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  50   resv2 += __nvvm_fmax_nan_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  51   resv2 += __nvvm_fmax_ftz_nan_f16x2(*(__fp16v2*)a , *(__fp16v2*)b);
  52   *out += __nvvm_fmax_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  53   *out += __nvvm_fmax_ftz_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  54   *out += __nvvm_fmax_nan_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  55   *out += __nvvm_fmax_ftz_nan_xorsign_abs_f16(*(__fp16*)a, *(__fp16*)b);
  56   resv2 += __nvvm_fmax_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  57   resv2 += __nvvm_fmax_ftz_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  58   resv2 += __nvvm_fmax_nan_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  59   resv2 += __nvvm_fmax_ftz_nan_xorsign_abs_f16x2(*(__fp16v2*)a, *(__fp16v2*)b);
  60
  61   *out += __nvvm_ldg_h((__fp16 *)a);
  62   resv2 += __nvvm_ldg_h2((__fp16v2 *)a);
  63
  64   *out += __nvvm_ldu_h((__fp16 *)a);
  65   resv2 += __nvvm_ldu_h2((__fp16v2 *)a);
  66
  67   *out += resv2[0] + resv2[1];
  68 }
  69
  70 // CHECK_ERROR: error: __nvvm_ex2_approx_f16 requires native half type support.
  71 // CHECK_ERROR: error: __nvvm_ex2_approx_f16x2 requires native half type support.
  72
  73 // CHECK_ERROR: error: __nvvm_fma_rn_relu_f16 requires native half type support.
  74 // CHECK_ERROR: error: __nvvm_fma_rn_ftz_relu_f16 requires native half type support.
  75 // CHECK_ERROR: error: __nvvm_fma_rn_relu_f16x2 requires native half type support.
  76 // CHECK_ERROR: error: __nvvm_fma_rn_ftz_relu_f16x2 requires native half type support.
  77 // CHECK_ERROR: error: __nvvm_fma_rn_ftz_f16 requires native half type support.
  78 // CHECK_ERROR: error: __nvvm_fma_rn_sat_f16 requires native half type support.
  79 // CHECK_ERROR: error: __nvvm_fma_rn_ftz_sat_f16 requires native half type support.
  80 // CHECK_ERROR: error: __nvvm_fma_rn_f16x2 requires native half type support.
  81 // CHECK_ERROR: error: __nvvm_fma_rn_ftz_f16x2 requires native half type support.
  82 // CHECK_ERROR: error: __nvvm_fma_rn_sat_f16x2 requires native half type support.
  83 // CHECK_ERROR: error: __nvvm_fma_rn_ftz_sat_f16x2 requires native half type support.
  84 // CHECK_ERROR: error: __nvvm_fmin_f16 requires native half type support.
  85 // CHECK_ERROR: error: __nvvm_fmin_ftz_f16 requires native half type support.
  86 // CHECK_ERROR: error: __nvvm_fmin_nan_f16 requires native half type support.
  87 // CHECK_ERROR: error: __nvvm_fmin_ftz_nan_f16 requires native half type support.
  88 // CHECK_ERROR: error: __nvvm_fmin_f16x2 requires native half type support.
  89 // CHECK_ERROR: error: __nvvm_fmin_ftz_f16x2 requires native half type support.
  90 // CHECK_ERROR: error: __nvvm_fmin_nan_f16x2 requires native half type support.
  91 // CHECK_ERROR: error: __nvvm_fmin_ftz_nan_f16x2 requires native half type support.
  92 // CHECK_ERROR: error: __nvvm_fmin_xorsign_abs_f16 requires native half type support.
  93 // CHECK_ERROR: error: __nvvm_fmin_ftz_xorsign_abs_f16 requires native half type support.
  94 // CHECK_ERROR: error: __nvvm_fmin_nan_xorsign_abs_f16 requires native half type support.
  95 // CHECK_ERROR: error: __nvvm_fmin_ftz_nan_xorsign_abs_f16 requires native half type support.
  96 // CHECK_ERROR: error: __nvvm_fmin_xorsign_abs_f16x2 requires native half type support.
  97 // CHECK_ERROR: error: __nvvm_fmin_ftz_xorsign_abs_f16x2 requires native half type support.
  98 // CHECK_ERROR: error: __nvvm_fmin_nan_xorsign_abs_f16x2 requires native half type support.
  99 // CHECK_ERROR: error: __nvvm_fmin_ftz_nan_xorsign_abs_f16x2 requires native half type support.
 100 // CHECK_ERROR: error: __nvvm_fmax_f16 requires native half type support.
 101 // CHECK_ERROR: error: __nvvm_fmax_ftz_f16 requires native half type support.
 102 // CHECK_ERROR: error: __nvvm_fmax_nan_f16 requires native half type support.
 103 // CHECK_ERROR: error: __nvvm_fmax_ftz_nan_f16 requires native half type support.
 104 // CHECK_ERROR: error: __nvvm_fmax_f16x2 requires native half type support.
 105 // CHECK_ERROR: error: __nvvm_fmax_ftz_f16x2 requires native half type support.
 106 // CHECK_ERROR: error: __nvvm_fmax_nan_f16x2 requires native half type support.
 107 // CHECK_ERROR: error: __nvvm_fmax_ftz_nan_f16x2 requires native half type support.
 108 // CHECK_ERROR: error: __nvvm_fmax_xorsign_abs_f16 requires native half type support.
 109 // CHECK_ERROR: error: __nvvm_fmax_ftz_xorsign_abs_f16 requires native half type support.
 110 // CHECK_ERROR: error: __nvvm_fmax_nan_xorsign_abs_f16 requires native half type support.
 111 // CHECK_ERROR: error: __nvvm_fmax_ftz_nan_xorsign_abs_f16 requires native half type support.
 112 // CHECK_ERROR: error: __nvvm_fmax_xorsign_abs_f16x2 requires native half type support.
 113 // CHECK_ERROR: error: __nvvm_fmax_ftz_xorsign_abs_f16x2 requires native half type support.
 114 // CHECK_ERROR: error: __nvvm_fmax_nan_xorsign_abs_f16x2 requires native half type support.
 115 // CHECK_ERROR: error: __nvvm_fmax_ftz_nan_xorsign_abs_f16x2 requires native half type support.
 116 // CHECK_ERROR: error: __nvvm_ldg_h requires native half type support.
 117 // CHECK_ERROR: error: __nvvm_ldg_h2 requires native half type support.
 118 // CHECK_ERROR: error: __nvvm_ldu_h requires native half type support.
 119 // CHECK_ERROR: error: __nvvm_ldu_h2 requires native half type support.