Skip to content

Instantly share code, notes, and snippets.

@etiennemlb
Last active May 25, 2025 17:16
Show Gist options
  • Save etiennemlb/4831ed004267e050ba96c1dbc248eaa6 to your computer and use it in GitHub Desktop.
Save etiennemlb/4831ed004267e050ba96c1dbc248eaa6 to your computer and use it in GitHub Desktop.
Clang builtins used in AMD HIP's CLR repo

Extracted from https://github.com/ROCm/clr/tree/7aae33951be2bfd9f25f0b9f89ee676b74d1a3bf

__builtin_amdgcn_alignbit
__builtin_amdgcn_atomic_dec32
__builtin_amdgcn_atomic_inc32
__builtin_amdgcn_ballot_w64
__builtin_amdgcn_cvt_f32_bf8
__builtin_amdgcn_cvt_f32_fp8
__builtin_amdgcn_cvt_pk_bf8_f32
__builtin_amdgcn_cvt_pk_f32_bf8
__builtin_amdgcn_cvt_pk_f32_fp8
__builtin_amdgcn_cvt_pk_fp8_f32
__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32
__builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32
__builtin_amdgcn_cvt_scalef32_f16_bf8
__builtin_amdgcn_cvt_scalef32_f16_fp8
__builtin_amdgcn_cvt_scalef32_f32_bf8
__builtin_amdgcn_cvt_scalef32_f32_fp8
__builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6
__builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6
__builtin_amdgcn_cvt_scalef32_pk32_bf6_bf16
__builtin_amdgcn_cvt_scalef32_pk32_bf6_f16
__builtin_amdgcn_cvt_scalef32_pk32_f16_bf6
__builtin_amdgcn_cvt_scalef32_pk32_f16_fp6
__builtin_amdgcn_cvt_scalef32_pk32_f32_bf6
__builtin_amdgcn_cvt_scalef32_pk32_f32_fp6
__builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16
__builtin_amdgcn_cvt_scalef32_pk32_fp6_f16
__builtin_amdgcn_cvt_scalef32_pk_bf16_bf8
__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4
__builtin_amdgcn_cvt_scalef32_pk_bf16_fp8
__builtin_amdgcn_cvt_scalef32_pk_bf8_bf16
__builtin_amdgcn_cvt_scalef32_pk_bf8_f16
__builtin_amdgcn_cvt_scalef32_pk_bf8_f32
__builtin_amdgcn_cvt_scalef32_pk_f16_bf8
__builtin_amdgcn_cvt_scalef32_pk_f16_fp4
__builtin_amdgcn_cvt_scalef32_pk_f16_fp8
__builtin_amdgcn_cvt_scalef32_pk_f32_bf8
__builtin_amdgcn_cvt_scalef32_pk_f32_fp4
__builtin_amdgcn_cvt_scalef32_pk_f32_fp8
__builtin_amdgcn_cvt_scalef32_pk_fp4_bf16
__builtin_amdgcn_cvt_scalef32_pk_fp4_f16
__builtin_amdgcn_cvt_scalef32_pk_fp4_f32
__builtin_amdgcn_cvt_scalef32_pk_fp8_bf16
__builtin_amdgcn_cvt_scalef32_pk_fp8_f16
__builtin_amdgcn_cvt_scalef32_pk_fp8_f32
__builtin_amdgcn_cvt_scalef32_sr_bf8_bf16
__builtin_amdgcn_cvt_scalef32_sr_bf8_f16
__builtin_amdgcn_cvt_scalef32_sr_bf8_f32
__builtin_amdgcn_cvt_scalef32_sr_fp8_bf16
__builtin_amdgcn_cvt_scalef32_sr_fp8_f16
__builtin_amdgcn_cvt_scalef32_sr_fp8_f32
__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f32
__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_bf16
__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f16
__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32
__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16
__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16
__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32
__builtin_amdgcn_cvt_sr_bf16_f32
__builtin_amdgcn_cvt_sr_bf8_f32
__builtin_amdgcn_cvt_sr_f16_f32
__builtin_amdgcn_cvt_sr_fp8_f32
__builtin_amdgcn_ds_atomic_fadd_f32
__builtin_amdgcn_ds_bpermute
__builtin_amdgcn_ds_permute
__builtin_amdgcn_ds_swizzle
__builtin_amdgcn_fence
__builtin_amdgcn_flat_atomic_fadd_f64
__builtin_amdgcn_flat_atomic_fadd_v2bf16
__builtin_amdgcn_flat_atomic_fadd_v2f16
__builtin_amdgcn_flat_atomic_fmax_f64
__builtin_amdgcn_flat_atomic_fmin_f64
__builtin_amdgcn_fmed3f
__builtin_amdgcn_global_atomic_fadd_f32
__builtin_amdgcn_groupstaticsize
__builtin_amdgcn_is_invocable
__builtin_amdgcn_is_private
__builtin_amdgcn_is_shared
__builtin_amdgcn_mbcnt_hi
__builtin_amdgcn_mbcnt_lo
__builtin_amdgcn_mov_dpp
__builtin_amdgcn_readfirstlane
__builtin_amdgcn_s_barrier
__builtin_amdgcn_s_getreg
__builtin_amdgcn_s_memtime
__builtin_amdgcn_wave_barrier
__builtin_amdgcn_wavefrontsize
__builtin_amdgcn_workgroup_id_x
__builtin_amdgcn_workitem_id_x
__builtin_bit_cast
__builtin_bitreverse32
__builtin_bitreverse64
__builtin_clz
__builtin_convertvector
__builtin_ctz
__builtin_ctzll
__builtin_elementwise_ceil
__builtin_elementwise_floor
__builtin_elementwise_fma
__builtin_elementwise_rint
__builtin_elementwise_trunc
__builtin_expect
__builtin_fmax
__builtin_fmin
__builtin_huge_val
__builtin_huge_valf
__builtin_memcpy
__builtin_nanf
__builtin_popcount
__builtin_popcountl
__builtin_popcountll
__builtin_readcyclecounter
__builtin_return_address
__builtin_trap


__ockl_activelane_u32
__ockl_atomic_add_noret_f32
__ockl_clz_u16
__ockl_clz_u32
__ockl_clz_u64
__ockl_clz_u8
__ockl_dm_init_v1
__ockl_fdot2
__ockl_fprintf_append_args
__ockl_fprintf_append_string_n
__ockl_fprintf_stderr_begin
__ockl_get_global_id
__ockl_get_group_id
__ockl_get_local_id
__ockl_get_local_size
__ockl_get_num_groups
__ockl_grid_is_valid
__ockl_grid_sync
__ockl_gws_barrier
__ockl_gws_init
__ockl_image_channel_data_type_1D
__ockl_image_channel_data_type_1Da
__ockl_image_channel_data_type_1Db
__ockl_image_channel_data_type_2D
__ockl_image_channel_data_type_2Da
__ockl_image_channel_data_type_2Dad
__ockl_image_channel_data_type_2Dd
__ockl_image_channel_data_type_3D
__ockl_image_channel_data_type_CM
__ockl_image_channel_data_type_CMa
__ockl_image_channel_order_1D
__ockl_image_channel_order_1Da
__ockl_image_channel_order_1Db
__ockl_image_channel_order_2D
__ockl_image_channel_order_2Da
__ockl_image_channel_order_2Dad
__ockl_image_channel_order_2Dd
__ockl_image_channel_order_3D
__ockl_image_channel_order_CM
__ockl_image_channel_order_CMa
__ockl_image_gather4a_2D
__ockl_image_gather4b_2D
__ockl_image_gather4g_2D
__ockl_image_gather4r_2D
__ockl_image_load_1D
__ockl_image_load_1Db
__ockl_image_load_2D
__ockl_image_load_3D
__ockl_image_load_CM
__ockl_image_load_lod_1D
__ockl_image_load_lod_2D
__ockl_image_load_lod_CM
__ockl_image_sample_1D
__ockl_image_sample_1Da
__ockl_image_sample_2D
__ockl_image_sample_2Da
__ockl_image_sample_3D
__ockl_image_sample_CM
__ockl_image_sample_CMa
__ockl_image_sample_grad_1D
__ockl_image_sample_grad_1Da
__ockl_image_sample_grad_2D
__ockl_image_sample_grad_2Da
__ockl_image_sample_grad_3D
__ockl_image_sample_grad_CM
__ockl_image_sample_grad_CMa
__ockl_image_sample_lod_1D
__ockl_image_sample_lod_1Da
__ockl_image_sample_lod_2D
__ockl_image_sample_lod_2Da
__ockl_image_sample_lod_3D
__ockl_image_sample_lod_CM
__ockl_image_sample_lod_CMa
__ockl_image_store_1D
__ockl_image_store_1Da
__ockl_image_store_2D
__ockl_image_store_2Da
__ockl_image_store_3D
__ockl_image_store_CM
__ockl_image_store_CMa
__ockl_image_store_lod_1D
__ockl_image_store_lod_1Da
__ockl_image_store_lod_2D
__ockl_image_store_lod_2Da
__ockl_image_store_lod_3D
__ockl_image_store_lod_CM
__ockl_image_store_lod_CMa
__ockl_lane_u32
__ockl_mul24_i32
__ockl_mul24_u32
__ockl_mul_hi_i32
__ockl_mul_hi_u32
__ockl_multi_grid_grid_rank
__ockl_multi_grid_is_valid
__ockl_multi_grid_num_grids
__ockl_multi_grid_size
__ockl_multi_grid_sync
__ockl_multi_grid_thread_rank
__ockl_sadd_u32
__ockl_sdot2
__ockl_sdot4
__ockl_sdot8
__ockl_steadyctr_u64
__ockl_udot2
__ockl_udot4
__ockl_udot8
__ockl_wfall_i32
__ockl_wfany_i32
__ockl_wfred_add_f16
__ockl_wfred_add_f32
__ockl_wfred_add_f64
__ockl_wfred_add_i32
__ockl_wfred_add_i64
__ockl_wfred_add_u32
__ockl_wfred_add_u64
__ockl_wfred_and_i32
__ockl_wfred_and_i64
__ockl_wfred_and_u32
__ockl_wfred_and_u64
__ockl_wfred_max_f16
__ockl_wfred_max_f32
__ockl_wfred_max_f64
__ockl_wfred_max_i32
__ockl_wfred_max_i64
__ockl_wfred_max_u32
__ockl_wfred_max_u64
__ockl_wfred_min_f16
__ockl_wfred_min_f32
__ockl_wfred_min_f64
__ockl_wfred_min_i32
__ockl_wfred_min_i64
__ockl_wfred_min_u32
__ockl_wfred_min_u64
__ockl_wfred_or_i32
__ockl_wfred_or_i64
__ockl_wfred_or_u32
__ockl_wfred_or_u64
__ockl_wfred_xor_i32
__ockl_wfred_xor_i64
__ockl_wfred_xor_u32
__ockl_wfred_xor_u64
__ockl_wgred_add_i32
__ockl_wgred_and_i32
__ockl_wgred_or_i32


__ocml_acos_f32
__ocml_acos_f64
__ocml_acosh_f32
__ocml_acosh_f64
__ocml_add_rte_f32
__ocml_add_rte_f64
__ocml_add_rtn_f32
__ocml_add_rtn_f64
__ocml_add_rtp_f32
__ocml_add_rtp_f64
__ocml_add_rtz_f32
__ocml_add_rtz_f64
__ocml_asin_f32
__ocml_asin_f64
__ocml_asinh_f32
__ocml_asinh_f64
__ocml_atan2_f32
__ocml_atan2_f64
__ocml_atan_f32
__ocml_atan_f64
__ocml_atanh_f32
__ocml_atanh_f64
__ocml_cbrt_f32
__ocml_cbrt_f64
__ocml_ceil_2f16
__ocml_ceil_f16
__ocml_ceil_f32
__ocml_ceil_f64
__ocml_copysign_f32
__ocml_copysign_f64
__ocml_cos_2f16
__ocml_cos_f16
__ocml_cos_f32
__ocml_cos_f64
__ocml_cosh_f32
__ocml_cosh_f64
__ocml_cospi_f32
__ocml_cospi_f64
__ocml_cvtrtn_f16_f32
__ocml_cvtrtn_f32_f64
__ocml_cvtrtn_f32_s32
__ocml_cvtrtn_f32_s64
__ocml_cvtrtn_f32_u32
__ocml_cvtrtn_f32_u64
__ocml_cvtrtn_f64_s64
__ocml_cvtrtn_f64_u64
__ocml_cvtrtp_f16_f32
__ocml_cvtrtp_f32_f64
__ocml_cvtrtp_f32_s32
__ocml_cvtrtp_f32_s64
__ocml_cvtrtp_f32_u32
__ocml_cvtrtp_f32_u64
__ocml_cvtrtp_f64_s64
__ocml_cvtrtp_f64_u64
__ocml_cvtrtz_f16_f32
__ocml_cvtrtz_f32_f64
__ocml_cvtrtz_f32_s32
__ocml_cvtrtz_f32_s64
__ocml_cvtrtz_f32_u32
__ocml_cvtrtz_f32_u64
__ocml_cvtrtz_f64_s64
__ocml_cvtrtz_f64_u64
__ocml_div_rte_f32
__ocml_div_rte_f64
__ocml_div_rtn_f32
__ocml_div_rtn_f64
__ocml_div_rtp_f32
__ocml_div_rtp_f64
__ocml_div_rtz_f32
__ocml_div_rtz_f64
__ocml_erfc_f32
__ocml_erfc_f64
__ocml_erfcinv_f32
__ocml_erfcinv_f64
__ocml_erfcx_f32
__ocml_erfcx_f64
__ocml_erf_f32
__ocml_erf_f64
__ocml_erfinv_f32
__ocml_erfinv_f64
__ocml_exp10_2f16
__ocml_exp10_f16
__ocml_exp10_f32
__ocml_exp10_f64
__ocml_exp2_2f16
__ocml_exp2_f16
__ocml_exp_2f16
__ocml_exp2_f32
__ocml_exp2_f64
__ocml_exp_f16
__ocml_exp_f32
__ocml_exp_f64
__ocml_expm1_f32
__ocml_expm1_f64
__ocml_fabs_2f16
__ocml_fabs_f16
__ocml_fabs_f32
__ocml_fabs_f64
__ocml_fdim_f32
__ocml_fdim_f64
__ocml_floor_2f16
__ocml_floor_f16
__ocml_floor_f32
__ocml_floor_f64
__ocml_fma_2f16
__ocml_fma_f16
__ocml_fma_f32
__ocml_fma_f64
__ocml_fma_rte_f32
__ocml_fma_rte_f64
__ocml_fma_rtn_f32
__ocml_fma_rtn_f64
__ocml_fma_rtp_f32
__ocml_fma_rtp_f64
__ocml_fma_rtz_f32
__ocml_fma_rtz_f64
__ocml_fmax_f16
__ocml_fmax_f32
__ocml_fmax_f64
__ocml_fmin_f16
__ocml_fmin_f32
__ocml_fmin_f64
__ocml_fmod_f32
__ocml_fmod_f64
__ocml_frexp_f32
__ocml_frexp_f64
__ocml_hypot_f32
__ocml_hypot_f64
__ocml_i0_f32
__ocml_i0_f64
__ocml_i1_f32
__ocml_i1_f64
__ocml_ilogb_f32
__ocml_ilogb_f64
__ocml_isfinite_f32
__ocml_isfinite_f64
__ocml_isinf_2f16
__ocml_isinf_f16
__ocml_isinf_f32
__ocml_isinf_f64
__ocml_isnan_2f16
__ocml_isnan_f16
__ocml_isnan_f32
__ocml_isnan_f64
__ocml_j0_f32
__ocml_j0_f64
__ocml_j1_f32
__ocml_j1_f64
__ocml_ldexp_f32
__ocml_ldexp_f64
__ocml_len3_f32
__ocml_len3_f64
__ocml_len4_f32
__ocml_len4_f64
__ocml_lgamma_f32
__ocml_lgamma_f64
__ocml_log10_2f16
__ocml_log10_f16
__ocml_log10_f32
__ocml_log10_f64
__ocml_log1p_f32
__ocml_log1p_f64
__ocml_log2_2f16
__ocml_log2_f16
__ocml_log_2f16
__ocml_log2_f32
__ocml_log2_f64
__ocml_logb_f32
__ocml_logb_f64
__ocml_log_f16
__ocml_log_f32
__ocml_log_f64
__ocml_modf_f32
__ocml_modf_f64
__ocml_mul_rte_f32
__ocml_mul_rte_f64
__ocml_mul_rtn_f32
__ocml_mul_rtn_f64
__ocml_mul_rtp_f32
__ocml_mul_rtp_f64
__ocml_mul_rtz_f32
__ocml_mul_rtz_f64
__ocml_native_cos_f32
__ocml_native_exp10_f32
__ocml_native_exp_f32
__ocml_native_log10_f32
__ocml_native_log2_f32
__ocml_native_log_f32
__ocml_native_sin_f32
__ocml_native_sqrt_f32
__ocml_ncdf_f32
__ocml_ncdf_f64
__ocml_ncdfinv_f32
__ocml_ncdfinv_f64
__ocml_nearbyint_f32
__ocml_nearbyint_f64
__ocml_nextafter_f32
__ocml_nextafter_f64
__ocml_pow_f32
__ocml_pow_f64
__ocml_pown_f16
__ocml_pown_f32
__ocml_pown_f64
__ocml_rcbrt_f32
__ocml_rcbrt_f64
__ocml_remainder_f32
__ocml_remainder_f64
__ocml_remquo_f32
__ocml_remquo_f64
__ocml_rhypot_f32
__ocml_rhypot_f64
__ocml_rint_2f16
__ocml_rint_f16
__ocml_rint_f32
__ocml_rint_f64
__ocml_rlen3_f32
__ocml_rlen3_f64
__ocml_rlen4_f32
__ocml_rlen4_f64
__ocml_round_f32
__ocml_round_f64
__ocml_rsqrt_2f16
__ocml_rsqrt_f16
__ocml_rsqrt_f32
__ocml_rsqrt_f64
__ocml_scalb_f32
__ocml_scalb_f64
__ocml_scalbn_f32
__ocml_scalbn_f64
__ocml_signbit_f32
__ocml_signbit_f64
__ocml_sin_2f16
__ocml_sincos_f32
__ocml_sincos_f64
__ocml_sincospi_f32
__ocml_sincospi_f64
__ocml_sin_f16
__ocml_sin_f32
__ocml_sin_f64
__ocml_sinh_f32
__ocml_sinh_f64
__ocml_sinpi_f32
__ocml_sinpi_f64
__ocml_sqrt_2f16
__ocml_sqrt_f16
__ocml_sqrt_f32
__ocml_sqrt_f64
__ocml_sqrt_rte_f32
__ocml_sqrt_rte_f64
__ocml_sqrt_rtn_f32
__ocml_sqrt_rtn_f64
__ocml_sqrt_rtp_f32
__ocml_sqrt_rtp_f64
__ocml_sqrt_rtz_f32
__ocml_sqrt_rtz_f64
__ocml_sub_rte_f32
__ocml_sub_rte_f64
__ocml_sub_rtn_f32
__ocml_sub_rtn_f64
__ocml_sub_rtp_f32
__ocml_sub_rtp_f64
__ocml_sub_rtz_f32
__ocml_sub_rtz_f64
__ocml_tan_f32
__ocml_tan_f64
__ocml_tanh_f32
__ocml_tanh_f64
__ocml_tgamma_f32
__ocml_tgamma_f64
__ocml_trunc_2f16
__ocml_trunc_f16
__ocml_trunc_f32
__ocml_trunc_f64
__ocml_y0_f32
__ocml_y0_f64
__ocml_y1_f32
__ocml_y1_f64
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment