diff options
author | Justin Holewinski <justin.holewinski@gmail.com> | 2012-11-09 23:50:51 +0000 |
---|---|---|
committer | Justin Holewinski <justin.holewinski@gmail.com> | 2012-11-09 23:50:51 +0000 |
commit | d1542c23b278f8420a938e8d6f192a55b225918e (patch) | |
tree | c9142d720be21186a104cf28fce580fee236321c | |
parent | da2d2c724da8a495533187c3c8d385004a7ec3d0 (diff) | |
download | clang-d1542c23b278f8420a938e8d6f192a55b225918e.tar.gz clang-d1542c23b278f8420a938e8d6f192a55b225918e.tar.bz2 clang-d1542c23b278f8420a938e8d6f192a55b225918e.tar.xz |
[NVPTX] Add __nvvm_* intrinsics as Clang builtins
Fixes bug 13354.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@167647 91177308-0d34-0410-b5e6-96231b3b80d8
-rw-r--r-- | include/clang/Basic/BuiltinsNVPTX.def | 246 | ||||
-rw-r--r-- | test/CodeGen/builtins-nvptx.c | 74 |
2 files changed, 318 insertions, 2 deletions
diff --git a/include/clang/Basic/BuiltinsNVPTX.def b/include/clang/Basic/BuiltinsNVPTX.def index f90a43f7f4..2fde59bdd7 100644 --- a/include/clang/Basic/BuiltinsNVPTX.def +++ b/include/clang/Basic/BuiltinsNVPTX.def @@ -14,6 +14,7 @@ // The format of this database matches clang/Basic/Builtins.def. +// Builtins retained from previous PTX back-end BUILTIN(__builtin_ptx_read_tid_x, "i", "nc") BUILTIN(__builtin_ptx_read_tid_y, "i", "nc") BUILTIN(__builtin_ptx_read_tid_z, "i", "nc") @@ -59,4 +60,249 @@ BUILTIN(__builtin_ptx_read_pm3, "i", "n") BUILTIN(__builtin_ptx_bar_sync, "vi", "n") +// Builtins exposed as part of NVVM +BUILTIN(__syncthreads, "v", "nc") +BUILTIN(__nvvm_bar0, "v", "nc") +BUILTIN(__nvvm_bar0_popc, "ii", "nc") +BUILTIN(__nvvm_bar0_and, "ii", "nc") +BUILTIN(__nvvm_bar0_or, "ii", "nc") +BUILTIN(__nvvm_membar_cta, "v", "n") +BUILTIN(__nvvm_membar_gl, "v", "n") +BUILTIN(__nvvm_membar_sys, "v", "n") +BUILTIN(__nvvm_popc_i, "ii", "n") +BUILTIN(__nvvm_popc_ll, "iLLi", "n") +BUILTIN(__nvvm_prmt, "UiUiUiUi", "nc") +BUILTIN(__nvvm_min_i, "iii", "nc") +BUILTIN(__nvvm_min_ui, "UiUiUi", "nc") +BUILTIN(__nvvm_min_ll, "LLiLLiLLi", "nc") +BUILTIN(__nvvm_min_ull, "ULLiULLiULLi", "nc") +BUILTIN(__nvvm_max_i, "iii", "nc") +BUILTIN(__nvvm_max_ui, "UiUiUi", "nc") +BUILTIN(__nvvm_max_ll, "LLiLLiLLi", "nc") +BUILTIN(__nvvm_max_ull, "ULLiULLiULLi", "nc") +BUILTIN(__nvvm_mulhi_i, "iii", "nc") +BUILTIN(__nvvm_mulhi_ui, "UiUiUi", "nc") +BUILTIN(__nvvm_mulhi_ll, "LLiLLiLLi", "nc") +BUILTIN(__nvvm_mulhi_ull, "ULLiULLiULLi", "nc") +BUILTIN(__nvvm_mul24_i, "iii", "nc") +BUILTIN(__nvvm_mul24_ui, "UiUiUi", "nc") +BUILTIN(__nvvm_brev32, "UiUi", "nc") +BUILTIN(__nvvm_brev64, "ULLiULLi", "nc") +BUILTIN(__nvvm_sad_i, "iiii", "nc") +BUILTIN(__nvvm_sad_ui, "UiUiUiUi", "nc") +BUILTIN(__nvvm_abs_i, "ii", "nc") +BUILTIN(__nvvm_abs_ll, "LiLi", "nc") +BUILTIN(__nvvm_floor_ftz_f, "ff", "nc") +BUILTIN(__nvvm_floor_f, "ff", "nc") +BUILTIN(__nvvm_floor_d, "dd", "nc") +BUILTIN(__nvvm_fabs_ftz_f, "ff", "nc") +BUILTIN(__nvvm_fabs_f, "ff", "nc") +BUILTIN(__nvvm_fabs_d, "dd", "nc") +BUILTIN(__nvvm_rcp_approx_ftz_d, "dd", "nc") +BUILTIN(__nvvm_fmin_ftz_f, "fff", "nc") +BUILTIN(__nvvm_fmin_f, "fff", "nc") +BUILTIN(__nvvm_fmax_ftz_f, "fff", "nc") +BUILTIN(__nvvm_fmax_f, "fff", "nc") +BUILTIN(__nvvm_rsqrt_approx_ftz_f, "ff", "nc") +BUILTIN(__nvvm_rsqrt_approx_f, "ff", "nc") +BUILTIN(__nvvm_fmin_d, "ddd", "nc") +BUILTIN(__nvvm_fmax_d, "ddd", "nc") +BUILTIN(__nvvm_rsqrt_approx_d, "dd", "nc") +BUILTIN(__nvvm_ceil_d, "dd", "nc") +BUILTIN(__nvvm_trunc_d, "dd", "nc") +BUILTIN(__nvvm_round_d, "dd", "nc") +BUILTIN(__nvvm_ex2_approx_d, "dd", "nc") +BUILTIN(__nvvm_lg2_approx_d, "dd", "nc") +BUILTIN(__nvvm_round_ftz_f, "ff", "nc") +BUILTIN(__nvvm_round_f, "ff", "nc") +BUILTIN(__nvvm_ex2_approx_ftz_f, "ff", "nc") +BUILTIN(__nvvm_ex2_approx_f, "ff", "nc") +BUILTIN(__nvvm_lg2_approx_ftz_f, "ff", "nc") +BUILTIN(__nvvm_lg2_approx_f, "ff", "nc") +BUILTIN(__nvvm_sin_approx_ftz_f, "ff", "nc") +BUILTIN(__nvvm_sin_approx_f, "ff", "nc") +BUILTIN(__nvvm_cos_approx_ftz_f, "ff", "nc") +BUILTIN(__nvvm_cos_approx_f, "ff", "nc") +BUILTIN(__nvvm_trunc_ftz_f, "ff", "nc") +BUILTIN(__nvvm_trunc_f, "ff", "nc") +BUILTIN(__nvvm_ceil_ftz_f, "ff", "nc") +BUILTIN(__nvvm_ceil_f, "ff", "nc") +BUILTIN(__nvvm_saturate_d, "dd", "nc") +BUILTIN(__nvvm_saturate_ftz_f, "ff", "nc") +BUILTIN(__nvvm_saturate_f, "ff", "nc") +BUILTIN(__nvvm_fma_rn_ftz_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rn_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rz_ftz_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rz_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rm_ftz_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rm_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rp_ftz_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rp_f, "ffff", "nc") +BUILTIN(__nvvm_fma_rn_d, "dddd", "nc") +BUILTIN(__nvvm_fma_rz_d, "dddd", "nc") +BUILTIN(__nvvm_fma_rm_d, "dddd", "nc") +BUILTIN(__nvvm_fma_rp_d, "dddd", "nc") +BUILTIN(__nvvm_div_approx_ftz_f, "fff", "nc") +BUILTIN(__nvvm_div_approx_f, "fff", "nc") +BUILTIN(__nvvm_div_rn_ftz_f, "fff", "nc") +BUILTIN(__nvvm_div_rn_f, "fff", "nc") +BUILTIN(__nvvm_div_rz_ftz_f, "fff", "nc") +BUILTIN(__nvvm_div_rz_f, "fff", "nc") +BUILTIN(__nvvm_div_rm_ftz_f, "fff", "nc") +BUILTIN(__nvvm_div_rm_f, "fff", "nc") +BUILTIN(__nvvm_div_rp_ftz_f, "fff", "nc") +BUILTIN(__nvvm_div_rp_f, "fff", "nc") +BUILTIN(__nvvm_rcp_rn_ftz_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rn_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rz_ftz_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rz_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rm_ftz_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rm_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rp_ftz_f, "ff", "nc") +BUILTIN(__nvvm_rcp_rp_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rn_ftz_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rn_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rz_ftz_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rz_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rm_ftz_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rm_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rp_ftz_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_rp_f, "ff", "nc") +BUILTIN(__nvvm_div_rn_d, "ddd", "nc") +BUILTIN(__nvvm_div_rz_d, "ddd", "nc") +BUILTIN(__nvvm_div_rm_d, "ddd", "nc") +BUILTIN(__nvvm_div_rp_d, "ddd", "nc") +BUILTIN(__nvvm_rcp_rn_d, "dd", "nc") +BUILTIN(__nvvm_rcp_rz_d, "dd", "nc") +BUILTIN(__nvvm_rcp_rm_d, "dd", "nc") +BUILTIN(__nvvm_rcp_rp_d, "dd", "nc") +BUILTIN(__nvvm_sqrt_rn_d, "dd", "nc") +BUILTIN(__nvvm_sqrt_rz_d, "dd", "nc") +BUILTIN(__nvvm_sqrt_rm_d, "dd", "nc") +BUILTIN(__nvvm_sqrt_rp_d, "dd", "nc") +BUILTIN(__nvvm_sqrt_approx_ftz_f, "ff", "nc") +BUILTIN(__nvvm_sqrt_approx_f, "ff", "nc") +BUILTIN(__nvvm_add_rn_d, "ddd", "nc") +BUILTIN(__nvvm_add_rz_d, "ddd", "nc") +BUILTIN(__nvvm_add_rm_d, "ddd", "nc") +BUILTIN(__nvvm_add_rp_d, "ddd", "nc") +BUILTIN(__nvvm_mul_rn_d, "ddd", "nc") +BUILTIN(__nvvm_mul_rz_d, "ddd", "nc") +BUILTIN(__nvvm_mul_rm_d, "ddd", "nc") +BUILTIN(__nvvm_mul_rp_d, "ddd", "nc") +BUILTIN(__nvvm_add_rm_ftz_f, "fff", "nc") +BUILTIN(__nvvm_add_rm_f, "fff", "nc") +BUILTIN(__nvvm_add_rp_ftz_f, "fff", "nc") +BUILTIN(__nvvm_add_rp_f, "fff", "nc") +BUILTIN(__nvvm_mul_rm_ftz_f, "fff", "nc") +BUILTIN(__nvvm_mul_rm_f, "fff", "nc") +BUILTIN(__nvvm_mul_rp_ftz_f, "fff", "nc") +BUILTIN(__nvvm_mul_rp_f, "fff", "nc") +BUILTIN(__nvvm_add_rn_ftz_f, "fff", "nc") +BUILTIN(__nvvm_add_rn_f, "fff", "nc") +BUILTIN(__nvvm_add_rz_ftz_f, "fff", "nc") +BUILTIN(__nvvm_add_rz_f, "fff", "nc") +BUILTIN(__nvvm_mul_rn_ftz_f, "fff", "nc") +BUILTIN(__nvvm_mul_rn_f, "fff", "nc") +BUILTIN(__nvvm_mul_rz_ftz_f, "fff", "nc") +BUILTIN(__nvvm_mul_rz_f, "fff", "nc") +BUILTIN(__nvvm_d2f_rn_ftz, "fd", "nc") +BUILTIN(__nvvm_d2f_rn, "fd", "nc") +BUILTIN(__nvvm_d2f_rz_ftz, "fd", "nc") +BUILTIN(__nvvm_d2f_rz, "fd", "nc") +BUILTIN(__nvvm_d2f_rm_ftz, "fd", "nc") +BUILTIN(__nvvm_d2f_rm, "fd", "nc") +BUILTIN(__nvvm_d2f_rp_ftz, "fd", "nc") +BUILTIN(__nvvm_d2f_rp, "fd", "nc") +BUILTIN(__nvvm_d2i_rn, "id", "nc") +BUILTIN(__nvvm_d2i_rz, "id", "nc") +BUILTIN(__nvvm_d2i_rm, "id", "nc") +BUILTIN(__nvvm_d2i_rp, "id", "nc") +BUILTIN(__nvvm_d2ui_rn, "Uid", "nc") +BUILTIN(__nvvm_d2ui_rz, "Uid", "nc") +BUILTIN(__nvvm_d2ui_rm, "Uid", "nc") +BUILTIN(__nvvm_d2ui_rp, "Uid", "nc") +BUILTIN(__nvvm_i2d_rn, "di", "nc") +BUILTIN(__nvvm_i2d_rz, "di", "nc") +BUILTIN(__nvvm_i2d_rm, "di", "nc") +BUILTIN(__nvvm_i2d_rp, "di", "nc") +BUILTIN(__nvvm_ui2d_rn, "dUi", "nc") +BUILTIN(__nvvm_ui2d_rz, "dUi", "nc") +BUILTIN(__nvvm_ui2d_rm, "dUi", "nc") +BUILTIN(__nvvm_ui2d_rp, "dUi", "nc") +BUILTIN(__nvvm_f2i_rn_ftz, "if", "nc") +BUILTIN(__nvvm_f2i_rn, "if", "nc") +BUILTIN(__nvvm_f2i_rz_ftz, "if", "nc") +BUILTIN(__nvvm_f2i_rz, "if", "nc") +BUILTIN(__nvvm_f2i_rm_ftz, "if", "nc") +BUILTIN(__nvvm_f2i_rm, "if", "nc") +BUILTIN(__nvvm_f2i_rp_ftz, "if", "nc") +BUILTIN(__nvvm_f2i_rp, "if", "nc") +BUILTIN(__nvvm_f2ui_rn_ftz, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rn, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rz_ftz, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rz, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rm_ftz, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rm, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rp_ftz, "Uif", "nc") +BUILTIN(__nvvm_f2ui_rp, "Uif", "nc") +BUILTIN(__nvvm_i2f_rn, "fi", "nc") +BUILTIN(__nvvm_i2f_rz, "fi", "nc") +BUILTIN(__nvvm_i2f_rm, "fi", "nc") +BUILTIN(__nvvm_i2f_rp, "fi", "nc") +BUILTIN(__nvvm_ui2f_rn, "fUi", "nc") +BUILTIN(__nvvm_ui2f_rz, "fUi", "nc") +BUILTIN(__nvvm_ui2f_rm, "fUi", "nc") +BUILTIN(__nvvm_ui2f_rp, "fUi", "nc") +BUILTIN(__nvvm_lohi_i2d, "dii", "nc") +BUILTIN(__nvvm_d2i_lo, "id", "nc") +BUILTIN(__nvvm_d2i_hi, "id", "nc") +BUILTIN(__nvvm_f2ll_rn_ftz, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rn, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rz_ftz, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rz, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rm_ftz, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rm, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rp_ftz, "LLif", "nc") +BUILTIN(__nvvm_f2ll_rp, "LLif", "nc") +BUILTIN(__nvvm_f2ull_rn_ftz, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rn, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rz_ftz, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rz, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rm_ftz, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rm, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rp_ftz, "ULLif", "nc") +BUILTIN(__nvvm_f2ull_rp, "ULLif", "nc") +BUILTIN(__nvvm_d2ll_rn, "LLid", "nc") +BUILTIN(__nvvm_d2ll_rz, "LLid", "nc") +BUILTIN(__nvvm_d2ll_rm, "LLid", "nc") +BUILTIN(__nvvm_d2ll_rp, "LLid", "nc") +BUILTIN(__nvvm_d2ull_rn, "ULLid", "nc") +BUILTIN(__nvvm_d2ull_rz, "ULLid", "nc") +BUILTIN(__nvvm_d2ull_rm, "ULLid", "nc") +BUILTIN(__nvvm_d2ull_rp, "ULLid", "nc") +BUILTIN(__nvvm_ll2f_rn, "fLLi", "nc") +BUILTIN(__nvvm_ll2f_rz, "fLLi", "nc") +BUILTIN(__nvvm_ll2f_rm, "fLLi", "nc") +BUILTIN(__nvvm_ll2f_rp, "fLLi", "nc") +BUILTIN(__nvvm_ull2f_rn, "fULLi", "nc") +BUILTIN(__nvvm_ull2f_rz, "fULLi", "nc") +BUILTIN(__nvvm_ull2f_rm, "fULLi", "nc") +BUILTIN(__nvvm_ull2f_rp, "fULLi", "nc") +BUILTIN(__nvvm_ll2d_rn, "dLLi", "nc") +BUILTIN(__nvvm_ll2d_rz, "dLLi", "nc") +BUILTIN(__nvvm_ll2d_rm, "dLLi", "nc") +BUILTIN(__nvvm_ll2d_rp, "dLLi", "nc") +BUILTIN(__nvvm_ull2d_rn, "dULLi", "nc") +BUILTIN(__nvvm_ull2d_rz, "dULLi", "nc") +BUILTIN(__nvvm_ull2d_rm, "dULLi", "nc") +BUILTIN(__nvvm_ull2d_rp, "dULLi", "nc") +BUILTIN(__nvvm_f2h_rn_ftz, "Usf", "nc") +BUILTIN(__nvvm_f2h_rn, "Usf", "nc") +BUILTIN(__nvvm_h2f, "fUs", "nc") +BUILTIN(__nvvm_bitcast_i2f, "fi", "nc") +BUILTIN(__nvvm_bitcast_f2i, "if", "nc") +BUILTIN(__nvvm_bitcast_ll2d, "dLLi", "nc") +BUILTIN(__nvvm_bitcast_d2ll, "LLid", "nc") + #undef BUILTIN diff --git a/test/CodeGen/builtins-nvptx.c b/test/CodeGen/builtins-nvptx.c index fa6b14c1ca..2c7e0c1367 100644 --- a/test/CodeGen/builtins-nvptx.c +++ b/test/CodeGen/builtins-nvptx.c @@ -1,8 +1,15 @@ -// RUN: %clang_cc1 -triple nvptx-unknown-unknown -emit-llvm -o %t %s -// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -emit-llvm -o %t %s +// REQUIRES: nvptx-registered-target +// REQUIRES: nvptx64-registered-target +// RUN: %clang_cc1 -triple nvptx-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s int read_tid() { +// CHECK: call i32 @llvm.ptx.read.tid.x() +// CHECK: call i32 @llvm.ptx.read.tid.y() +// CHECK: call i32 @llvm.ptx.read.tid.z() +// CHECK: call i32 @llvm.ptx.read.tid.w() + int x = __builtin_ptx_read_tid_x(); int y = __builtin_ptx_read_tid_y(); int z = __builtin_ptx_read_tid_z(); @@ -14,6 +21,11 @@ int read_tid() { int read_ntid() { +// CHECK: call i32 @llvm.ptx.read.ntid.x() +// CHECK: call i32 @llvm.ptx.read.ntid.y() +// CHECK: call i32 @llvm.ptx.read.ntid.z() +// CHECK: call i32 @llvm.ptx.read.ntid.w() + int x = __builtin_ptx_read_ntid_x(); int y = __builtin_ptx_read_ntid_y(); int z = __builtin_ptx_read_ntid_z(); @@ -25,6 +37,11 @@ int read_ntid() { int read_ctaid() { +// CHECK: call i32 @llvm.ptx.read.ctaid.x() +// CHECK: call i32 @llvm.ptx.read.ctaid.y() +// CHECK: call i32 @llvm.ptx.read.ctaid.z() +// CHECK: call i32 @llvm.ptx.read.ctaid.w() + int x = __builtin_ptx_read_ctaid_x(); int y = __builtin_ptx_read_ctaid_y(); int z = __builtin_ptx_read_ctaid_z(); @@ -36,6 +53,11 @@ int read_ctaid() { int read_nctaid() { +// CHECK: call i32 @llvm.ptx.read.nctaid.x() +// CHECK: call i32 @llvm.ptx.read.nctaid.y() +// CHECK: call i32 @llvm.ptx.read.nctaid.z() +// CHECK: call i32 @llvm.ptx.read.nctaid.w() + int x = __builtin_ptx_read_nctaid_x(); int y = __builtin_ptx_read_nctaid_y(); int z = __builtin_ptx_read_nctaid_z(); @@ -47,6 +69,13 @@ int read_nctaid() { int read_ids() { +// CHECK: call i32 @llvm.ptx.read.laneid() +// CHECK: call i32 @llvm.ptx.read.warpid() +// CHECK: call i32 @llvm.ptx.read.nwarpid() +// CHECK: call i32 @llvm.ptx.read.smid() +// CHECK: call i32 @llvm.ptx.read.nsmid() +// CHECK: call i32 @llvm.ptx.read.gridid() + int a = __builtin_ptx_read_laneid(); int b = __builtin_ptx_read_warpid(); int c = __builtin_ptx_read_nwarpid(); @@ -60,6 +89,12 @@ int read_ids() { int read_lanemasks() { +// CHECK: call i32 @llvm.ptx.read.lanemask.eq() +// CHECK: call i32 @llvm.ptx.read.lanemask.le() +// CHECK: call i32 @llvm.ptx.read.lanemask.lt() +// CHECK: call i32 @llvm.ptx.read.lanemask.ge() +// CHECK: call i32 @llvm.ptx.read.lanemask.gt() + int a = __builtin_ptx_read_lanemask_eq(); int b = __builtin_ptx_read_lanemask_le(); int c = __builtin_ptx_read_lanemask_lt(); @@ -73,6 +108,9 @@ int read_lanemasks() { long read_clocks() { +// CHECK: call i32 @llvm.ptx.read.clock() +// CHECK: call i64 @llvm.ptx.read.clock64() + int a = __builtin_ptx_read_clock(); long b = __builtin_ptx_read_clock64(); @@ -82,6 +120,11 @@ long read_clocks() { int read_pms() { +// CHECK: call i32 @llvm.ptx.read.pm0() +// CHECK: call i32 @llvm.ptx.read.pm1() +// CHECK: call i32 @llvm.ptx.read.pm2() +// CHECK: call i32 @llvm.ptx.read.pm3() + int a = __builtin_ptx_read_pm0(); int b = __builtin_ptx_read_pm1(); int c = __builtin_ptx_read_pm2(); @@ -93,6 +136,33 @@ int read_pms() { void sync() { +// CHECK: call void @llvm.ptx.bar.sync(i32 0) + __builtin_ptx_bar_sync(0); } + + +// NVVM intrinsics + +// The idea is not to test all intrinsics, just that Clang is recognizing the +// builtins defined in BuiltinsNVPTX.def +void nvvm_math(float f1, float f2, double d1, double d2) { +// CHECK: call float @llvm.nvvm.fmax.f + float t1 = __nvvm_fmax_f(f1, f2); +// CHECK: call float @llvm.nvvm.fmin.f + float t2 = __nvvm_fmin_f(f1, f2); +// CHECK: call float @llvm.nvvm.sqrt.rn.f + float t3 = __nvvm_sqrt_rn_f(f1); +// CHECK: call float @llvm.nvvm.rcp.rn.f + float t4 = __nvvm_rcp_rn_f(f2); + +// CHECK: call double @llvm.nvvm.fmax.d + double td1 = __nvvm_fmax_d(d1, d2); +// CHECK: call double @llvm.nvvm.fmin.d + double td2 = __nvvm_fmin_d(d1, d2); +// CHECK: call double @llvm.nvvm.sqrt.rn.d + double td3 = __nvvm_sqrt_rn_d(d1); +// CHECK: call double @llvm.nvvm.rcp.rn.d + double td4 = __nvvm_rcp_rn_d(d2); +} |