[NVPTX] Add __nvvm_* intrinsics as Clang builtins
Fixes bug 13354. llvm-svn: 167647
This commit is contained in:
Justin Holewinski
parent
2dc9d072e5
commit
c0cad046b6
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue