diff --git a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp index c59c69ffd9..dfce53bd26 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp @@ -35,7 +35,7 @@ THE SOFTWARE. #include #include -namespace glo_tests +namespace hip_impl { namespace { @@ -54,7 +54,6 @@ namespace glo_tests template requires(Domain == {Ts...}) - static inline void grid_launch_impl( New_grid_launch_tag, @@ -79,14 +78,13 @@ namespace glo_tests throw std::runtime_error{"Failed to retrieve accelerator_view!"}; } - hc::parallel_for_each(*av, d, [=](hc::tiled_index<3> idx) [[hc]] { + hc::parallel_for_each(*av, d, [=](const hc::tiled_index<3>& idx) [[hc]] { k(args...); }); } template requires(Domain == {hipLaunchParm, Ts...}) - static inline void grid_launch_impl( Old_grid_launch_tag, @@ -110,7 +108,6 @@ namespace glo_tests template requires(Domain == {Ts...}) - static inline std::enable_if_t::value> grid_launch( dim3 num_blocks, @@ -152,7 +149,6 @@ namespace glo_tests template requires(Domain == {Ts...}) - static inline void grid_launch( New_grid_launch_tag, @@ -174,7 +170,6 @@ namespace glo_tests template requires(Domain == {Ts...}) - static inline void grid_launch( Old_grid_launch_tag, @@ -196,7 +191,6 @@ namespace glo_tests template requires(Domain == {Ts...}) - static inline std::enable_if_t::value> grid_launch( dim3 num_blocks, @@ -214,402 +208,363 @@ namespace glo_tests std::forward(args)...); } - namespace - { - template struct Wrapper; - - template - struct Wrapper::value>> { - template - requires(Domain == {Ts...}) - void operator()(Ts&&... args) const - { - grid_launch(std::forward(args)...); - } - }; - - template - struct Wrapper::value>> { - template - void operator()(Ts&&...) const {} - }; - } - - // TODO: these are temporary, they need to be uglified and them completely - // removed once we enable C++14 support and can have proper generic, - // variadic lambdas. - #define make_lambda_wrapper21( \ - kernel_name, \ - p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \ - p16, p17, p18, p19) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_, \ - std::decay_t _p14_, \ - std::decay_t _p15_, \ - std::decay_t _p16_, \ - std::decay_t _p17_, \ - std::decay_t _p18_, \ - std::decay_t _p19_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_, \ - _p18_, _p19_); \ + // TODO: these are temporary, they need to be completely removed once we + // enable C++14 support and can have proper generic, variadic lambdas. + #define make_kernel_lambda_hip_21(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18, p19)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_,\ + const std::decay_t& _p19_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_, _p19_);\ } - #define make_lambda_wrapper20( \ - kernel_name, \ - p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \ - p16, p17, p18) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_, \ - std::decay_t _p14_, \ - std::decay_t _p15_, \ - std::decay_t _p16_, \ - std::decay_t _p17_, \ - std::decay_t _p18_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_, \ - _p18_); \ + #define make_kernel_lambda_hip_20(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17, p18)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_,\ + const std::decay_t& _p18_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_,\ + _p18_);\ } - #define make_lambda_wrapper19( \ - kernel_name, \ - p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \ - p16, p17) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_, \ - std::decay_t _p14_, \ - std::decay_t _p15_, \ - std::decay_t _p16_, \ - std::decay_t _p17_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_); \ + #define make_kernel_lambda_hip_19(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16, p17)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_,\ + const std::decay_t& _p17_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_, _p17_);\ } - #define make_lambda_wrapper18( \ - kernel_name, \ - p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, \ - p16) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_, \ - std::decay_t _p14_, \ - std::decay_t _p15_, \ - std::decay_t _p16_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_); \ + #define make_kernel_lambda_hip_18(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15,\ + p16)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_,\ + const std::decay_t& _p16_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_, _p16_);\ } - #define make_lambda_wrapper17( \ - kernel_name, \ - p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_, \ - std::decay_t _p14_, \ - std::decay_t _p15_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_); \ + #define make_kernel_lambda_hip_17(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_,\ + const std::decay_t& _p15_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_, _p15_);\ } - #define make_lambda_wrapper16( \ - kernel_name, \ - p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_, \ - std::decay_t _p14_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_, _p14_); \ + #define make_kernel_lambda_hip_16(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_,\ + const std::decay_t& _p14_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_, _p14_);\ } - #define make_lambda_wrapper15( \ - kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13)\ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_, \ - std::decay_t _p13_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_, _p13_); \ + #define make_kernel_lambda_hip_15(\ + kernel_name,\ + p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_,\ + const std::decay_t& _p13_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_, _p13_);\ } - #define make_lambda_wrapper14( \ - kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_, \ - std::decay_t _p12_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_, _p12_); \ + #define make_kernel_lambda_hip_14(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_,\ + const std::decay_t& _p12_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_, _p12_);\ } - #define make_lambda_wrapper13( \ - kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_, \ - std::decay_t _p11_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, \ - _p9_, _p10_, _p11_); \ + #define make_kernel_lambda_hip_13(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_,\ + const std::decay_t& _p11_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_,\ + _p9_, _p10_, _p11_);\ } - #define make_lambda_wrapper12( \ - kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_, \ - std::decay_t _p10_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_, \ - _p10_); \ + #define make_kernel_lambda_hip_12(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_,\ + const std::decay_t& _p10_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_,\ + _p10_);\ } - #define make_lambda_wrapper11( \ - kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_, \ - std::decay_t _p9_) [[hc]] { \ - kernel_name( \ - _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_); \ + #define make_kernel_lambda_hip_11(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8, p9)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_,\ + const std::decay_t& _p9_) [[hc]] {\ + kernel_name(\ + _p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_, _p9_);\ } - #define make_lambda_wrapper10( \ - kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_, \ - std::decay_t _p8_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_); \ + #define make_kernel_lambda_hip_10(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7, p8)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_,\ + const std::decay_t& _p8_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_, _p8_);\ } - #define make_lambda_wrapper9(kernel_name, p0, p1, p2, p3, p4, p5, p6, p7) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_, \ - std::decay_t _p7_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_); \ + #define make_kernel_lambda_hip_9(\ + kernel_name, p0, p1, p2, p3, p4, p5, p6, p7)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_,\ + const std::decay_t& _p7_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_, _p7_);\ } - #define make_lambda_wrapper8(kernel_name, p0, p1, p2, p3, p4, p5, p6) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_, \ - std::decay_t _p6_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_); \ + #define make_kernel_lambda_hip_8(kernel_name, p0, p1, p2, p3, p4, p5, p6)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_,\ + const std::decay_t& _p6_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_, _p6_);\ } - #define make_lambda_wrapper7(kernel_name, p0, p1, p2, p3, p4, p5) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_, \ - std::decay_t _p5_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_); \ + #define make_kernel_lambda_hip_7(kernel_name, p0, p1, p2, p3, p4, p5)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_,\ + const std::decay_t& _p5_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_, _p5_);\ } - #define make_lambda_wrapper6(kernel_name, p0, p1, p2, p3, p4) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_, \ - std::decay_t _p4_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_); \ + #define make_kernel_lambda_hip_6(kernel_name, p0, p1, p2, p3, p4)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_,\ + const std::decay_t& _p4_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_, _p4_);\ } - #define make_lambda_wrapper5(kernel_name, p0, p1, p2, p3) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_, \ - std::decay_t _p3_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_, _p3_); \ + #define make_kernel_lambda_hip_5(kernel_name, p0, p1, p2, p3)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_,\ + const std::decay_t& _p3_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_, _p3_);\ } - #define make_lambda_wrapper4(kernel_name, p0, p1, p2) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_, \ - std::decay_t _p2_) [[hc]] { \ - kernel_name(_p0_, _p1_, _p2_); \ + #define make_kernel_lambda_hip_4(kernel_name, p0, p1, p2)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_,\ + const std::decay_t& _p2_) [[hc]] {\ + kernel_name(_p0_, _p1_, _p2_);\ } - #define make_lambda_wrapper3(kernel_name, p0, p1) \ - []( \ - std::decay_t _p0_, \ - std::decay_t _p1_) [[hc]] { \ - kernel_name(_p0_, _p1_); \ + #define make_kernel_lambda_hip_3(kernel_name, p0, p1)\ + [](const std::decay_t& _p0_,\ + const std::decay_t& _p1_) [[hc]] {\ + kernel_name(_p0_, _p1_);\ } - #define make_lambda_wrapper2(kernel_name, p0) \ - [](std::decay_t _p0_) [[hc]] { \ - kernel_name(_p0_); \ + #define make_kernel_lambda_hip_2(kernel_name, p0)\ + [](const std::decay_t& _p0_) [[hc]] {\ + kernel_name(_p0_);\ } - #define make_lambda_wrapper1(kernel_name) \ + #define make_kernel_lambda_hip_1(kernel_name)\ []() [[hc]] { kernel_name(lp); } - #define make_lambda_wrapper(...) \ - overload_macro(make_lambda_wrapper, __VA_ARGS__) + #define make_kernel_lambda_hip_(...)\ + overload_macro_hip_(make_kernel_lambda_hip_, __VA_ARGS__) - #define hipLaunchKernelV3( \ - kernel_name, \ - num_blocks, \ - dim_blocks, \ - group_mem_bytes, \ - stream, \ - ...) \ - { \ - glo_tests::grid_launch( \ - num_blocks, \ - dim_blocks, \ - group_mem_bytes, \ - stream, \ - make_lambda_wrapper(kernel_name, __VA_ARGS__), \ - ##__VA_ARGS__); \ + #define hipLaunchKernelV3(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + ...)\ + {\ + hip_impl::grid_launch(\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + make_kernel_lambda_hip_(kernel_name, __VA_ARGS__),\ + ##__VA_ARGS__);\ } - #define hipLaunchKernel( \ - kernel_name, \ - num_blocks, \ - dim_blocks, \ - group_mem_bytes, \ - stream, \ - ...) \ - { \ - hipLaunchKernelV3( \ - kernel_name, \ - num_blocks, \ - dim_blocks, \ - group_mem_bytes, \ - stream, \ - hipLaunchParm{}, \ - ##__VA_ARGS__); \ + #define hipLaunchKernel(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + ...)\ + {\ + hipLaunchKernelV3(\ + kernel_name,\ + num_blocks,\ + dim_blocks,\ + group_mem_bytes,\ + stream,\ + hipLaunchParm{},\ + ##__VA_ARGS__);\ } } diff --git a/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp b/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp index 301d740066..566612dce4 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp @@ -20,9 +20,6 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -// -// Created by alexv on 08/11/16. -// #pragma once #include // For std::conditional, std::decay, std::enable_if, @@ -30,7 +27,7 @@ THE SOFTWARE. #include // For std::declval. namespace std -{ +{ // TODO: these should be removed as soon as possible. #if (__cplusplus < 201406L) template using void_t = void; @@ -126,20 +123,20 @@ namespace std #endif } -namespace glo_tests // Only for documentation, macros ignore namespaces. +namespace // Only for documentation, macros ignore namespaces. { - #define count_macro_args_impl( \ - _0, _1, _2, _3, _4, _5, _6, _7, \ - _8, _9, _10, _11, _12, _13, _14, _15, \ - _16, _17, _18, _19, _20, _21, _n, ...) _n - #define count_macro_args(...) \ - count_macro_args_impl( \ + #define count_macro_args_impl_hip_(\ + _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15,\ + _16, _17, _18, _19, _20, _21, _n, ...) _n + #define count_macro_args_hip_(...)\ + count_macro_args_impl_hip_(\ , ##__VA_ARGS__, 21, 20, 19, 18, 17, 16, 15, 14, 13, 12, 11, 10, 9,\ 8, 7, 6, 5, 4, 3, 2, 1, 0) - #define overloaded_macro_expand(macro, arg_cnt) macro##arg_cnt - #define overload_macro_impl(macro, arg_cnt) \ - overloaded_macro_expand(macro, arg_cnt) - #define overload_macro(macro, ...) \ - overload_macro_impl(macro, count_macro_args(__VA_ARGS__)) (__VA_ARGS__) + #define overloaded_macro_expand_hip_(macro, arg_cnt) macro##arg_cnt + #define overload_macro_impl_hip_(macro, arg_cnt)\ + overloaded_macro_expand_hip_(macro, arg_cnt) + #define overload_macro_hip_(macro, ...)\ + overload_macro_impl_hip_(macro, count_macro_args_hip_(__VA_ARGS__))\ + (__VA_ARGS__) } diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 590cc33bd1..42d80ecb0e 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -62,11 +62,15 @@ THE SOFTWARE. #include //TODO-HCC-GL - change this to typedef. //typedef grid_launch_parm hipLaunchParm ; -struct EmptyLaunchParm{}; + #if GENERIC_GRID_LAUNCH == 0 #define hipLaunchParm grid_launch_parm #else - #define hipLaunchParm EmptyLaunchParm +namespace hip_impl +{ + struct Empty_launch_parm{}; +} +#define hipLaunchParm hip_impl::Empty_launch_parm #endif //GENERIC_GRID_LAUNCH #if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20) || GENERIC_GRID_LAUNCH == 1 diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index cba6f01d49..658e3d2f97 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -32,7 +32,7 @@ THE SOFTWARE. #include #ifndef GENERIC_GRID_LAUNCH -#define GENERIC_GRID_LAUNCH 0 +#define GENERIC_GRID_LAUNCH 0 #endif #include diff --git a/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h b/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h index 28c9268f59..ab975dbb2c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h @@ -32,7 +32,7 @@ THE SOFTWARE. // Add guard to Generic Grid Launch method #ifndef GENERIC_GRID_LAUNCH -#define GENERIC_GRID_LAUNCH 0 +#define GENERIC_GRID_LAUNCH 0 #endif #ifdef __HCC__ @@ -46,7 +46,7 @@ THE SOFTWARE. #if GENERIC_GRID_LAUNCH == 0 //#warning "original global define reached" #define __global__ __attribute__((hc_grid_launch)) __attribute__((used)) -#else +#else //#warning "GGL global define reached" #define __global__ [[hc]] __attribute__((weak)) #endif //GENERIC_GRID_LAUNCH