210 lines · c
1// REQUIRES: nvptx-registered-target2//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 %s7 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 %s12 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 %s17 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 %s22 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 %s27 28// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \29// RUN: sm_53 -target-feature +ptx65 -fcuda-is-device -fnative-half-type \30// RUN: -emit-llvm -o - -x cuda %s \31// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX65_SM53 %s32 33// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown \34// RUN: -target-cpu sm_53 -target-feature +ptx65 -fcuda-is-device \35// RUN: -fnative-half-type -emit-llvm -o - -x cuda %s \36// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX65_SM53 %s37 38#define __device__ __attribute__((device))39 40__device__ void nvvm_ex2_sm75() {41#if __CUDA_ARCH__ >= 75042 // CHECK_PTX70_SM75: call half @llvm.nvvm.ex2.approx.f1643 __nvvm_ex2_approx_f16(0.1f16);44 // CHECK_PTX70_SM75: call <2 x half> @llvm.nvvm.ex2.approx.v2f1645 __nvvm_ex2_approx_f16x2({0.1f16, 0.7f16});46#endif47 // CHECK: ret void48}49 50// CHECK-LABEL: nvvm_min_max_sm8051__device__ void nvvm_min_max_sm80() {52#if __CUDA_ARCH__ >= 80053 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.f1654 __nvvm_fmin_f16(0.1f16, 0.1f16);55 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.ftz.f1656 __nvvm_fmin_ftz_f16(0.1f16, 0.1f16);57 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.nan.f1658 __nvvm_fmin_nan_f16(0.1f16, 0.1f16);59 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmin.ftz.nan.f1660 __nvvm_fmin_ftz_nan_f16(0.1f16, 0.1f16);61 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.f16x262 __nvvm_fmin_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});63 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.ftz.f16x264 __nvvm_fmin_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});65 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmin.nan.f16x266 __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.f16x268 __nvvm_fmin_ftz_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});69 70 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.f1671 __nvvm_fmax_f16(0.1f16, 0.1f16);72 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.ftz.f1673 __nvvm_fmax_ftz_f16(0.1f16, 0.1f16);74 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.nan.f1675 __nvvm_fmax_nan_f16(0.1f16, 0.1f16);76 // CHECK_PTX70_SM80: call half @llvm.nvvm.fmax.ftz.nan.f1677 __nvvm_fmax_ftz_nan_f16(0.1f16, 0.1f16);78 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.f16x279 __nvvm_fmax_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});80 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.ftz.f16x281 __nvvm_fmax_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});82 // CHECK_PTX70_SM80: call <2 x half> @llvm.nvvm.fmax.nan.f16x283 __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.f16x285 __nvvm_fmax_ftz_nan_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});86#endif87 // CHECK: ret void88}89 90// CHECK-LABEL: nvvm_fma_f16_f16x2_sm8091__device__ void nvvm_fma_f16_f16x2_sm80() {92#if __CUDA_ARCH__ >= 80093 // CHECK_PTX70_SM80: call half @llvm.nvvm.fma.rn.relu.f1694 __nvvm_fma_rn_relu_f16(0.1f16, 0.1f16, 0.1f16);95 // CHECK_PTX70_SM80: call half @llvm.nvvm.fma.rn.ftz.relu.f1696 __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.f16x299 __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.f16x2102 __nvvm_fma_rn_ftz_relu_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},103 {0.1f16, 0.7f16});104#endif105 // CHECK: ret void106}107 108// CHECK-LABEL: nvvm_fma_f16_f16x2_sm53109__device__ void nvvm_fma_f16_f16x2_sm53() {110#if __CUDA_ARCH__ >= 530111 // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.f16112 __nvvm_fma_rn_f16(0.1f16, 0.1f16, 0.1f16);113 // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.ftz.f16114 __nvvm_fma_rn_ftz_f16(0.1f16, 0.1f16, 0.1f16);115 // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.sat.f16116 __nvvm_fma_rn_sat_f16(0.1f16, 0.1f16, 0.1f16);117 // CHECK_PTX65_SM53: call half @llvm.nvvm.fma.rn.ftz.sat.f16118 __nvvm_fma_rn_ftz_sat_f16(0.1f16, 0.1f16, 0.1f16);119 120 // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.f16x2121 __nvvm_fma_rn_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},122 {0.1f16, 0.7f16});123 // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.f16x2124 __nvvm_fma_rn_ftz_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},125 {0.1f16, 0.7f16});126 // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.sat.f16x2127 __nvvm_fma_rn_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},128 {0.1f16, 0.7f16});129 // CHECK_PTX65_SM53: call <2 x half> @llvm.nvvm.fma.rn.ftz.sat.f16x2130 __nvvm_fma_rn_ftz_sat_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16},131 {0.1f16, 0.7f16});132#endif133 // CHECK: ret void134}135 136// CHECK-LABEL: nvvm_min_max_sm86137__device__ void nvvm_min_max_sm86() {138#if __CUDA_ARCH__ >= 860139 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.xorsign.abs.f16140 __nvvm_fmin_xorsign_abs_f16(0.1f16, 0.1f16);141 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.ftz.xorsign.abs.f16142 __nvvm_fmin_ftz_xorsign_abs_f16(0.1f16, 0.1f16);143 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.nan.xorsign.abs.f16144 __nvvm_fmin_nan_xorsign_abs_f16(0.1f16, 0.1f16);145 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmin.ftz.nan.xorsign.abs.f16146 __nvvm_fmin_ftz_nan_xorsign_abs_f16(0.1f16, 0.1f16);147 // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmin.xorsign.abs.f16x2148 __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.f16x2150 __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.f16x2152 __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.f16x2154 __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.f16157 __nvvm_fmax_xorsign_abs_f16(0.1f16, 0.1f16);158 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.ftz.xorsign.abs.f16159 __nvvm_fmax_ftz_xorsign_abs_f16(0.1f16, 0.1f16);160 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.nan.xorsign.abs.f16161 __nvvm_fmax_nan_xorsign_abs_f16(0.1f16, 0.1f16);162 // CHECK_PTX72_SM86: call half @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16163 __nvvm_fmax_ftz_nan_xorsign_abs_f16(0.1f16, 0.1f16);164 // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.xorsign.abs.f16x2165 __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.f16x2167 __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.f16x2169 __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.f16x2171 __nvvm_fmax_ftz_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});172#endif173 // CHECK: ret void174}175 176// CHECK-LABEL: nvvm_fabs_f16177__device__ void nvvm_fabs_f16() {178#if __CUDA_ARCH__ >= 530179 // CHECK: call half @llvm.nvvm.fabs.f16180 __nvvm_fabs_f16(0.1f16);181 // CHECK: call half @llvm.nvvm.fabs.ftz.f16182 __nvvm_fabs_ftz_f16(0.1f16);183 // CHECK: call <2 x half> @llvm.nvvm.fabs.v2f16184 __nvvm_fabs_f16x2({0.1f16, 0.7f16});185 // CHECK: call <2 x half> @llvm.nvvm.fabs.ftz.v2f16186 __nvvm_fabs_ftz_f16x2({0.1f16, 0.7f16});187#endif188 // CHECK: ret void189}190 191 192 193typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2)));194 195// CHECK-LABEL: nvvm_ldg_native_half_types196__device__ void nvvm_ldg_native_half_types(const void *p) {197 // CHECK: load half, ptr addrspace(1) {{.*}}, align 2, !invariant.load198 __nvvm_ldg_h((const __fp16 *)p);199 // CHECK: load <2 x half>, ptr addrspace(1) {{.*}}, align 4, !invariant.load200 __nvvm_ldg_h2((const __fp16v2 *)p);201}202 203// CHECK-LABEL: nvvm_ldu_native_half_types204__device__ void nvvm_ldu_native_half_types(const void *p) {205 // CHECK: call half @llvm.nvvm.ldu.global.f.f16.p0206 __nvvm_ldu_h((const __fp16 *)p);207 // CHECK: call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p0208 __nvvm_ldu_h2((const __fp16v2 *)p);209}210