Commit 833a31bb authored by Po-Yen, Chen's avatar Po-Yen, Chen
Browse files

Add macros to control IGLP builtin function usages

parent 3502b44e
......@@ -134,6 +134,8 @@
#define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 1
// experimental feature: add instances using pipeline v2
#define CK_EXPERIMENTAL_PIPELINE_V2_INSTANCES 1
// experimental feature: add __builtin_amdgcn_iglp_opt() optimized (pipeline) instances
#define CK_EXPERIMENTAL_IGLP_OPT_INSTANCES 0
// hack: have underlying assumption that need to be satsified, otherwise it's a bug
// hack for forcing register to keep idx_diff_low_const in SGPR. idx_diff_low_const must be
......@@ -161,6 +163,10 @@
#define CK_WORKAROUND_SWDEV_XXXXXX_BF16_ATTEN_FWD_GFX908_ISSUE 0
#endif // __gfx908__
// workaround: using __builtin_amdgcn_sched_barrier() or __builtin_amdgcn_sched_group_barrier() or
// __builtin_amdgcn_iglp_opt() may cause data hazard issues
#define CK_WORKAROUND_SWDEV_XXXXXX_SCHED_GROUP_DATA_HAZARD_ISSUE 1
namespace ck {
enum struct InMemoryDataOperationEnum
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment