diff options
Diffstat (limited to 'sci-libs/miopen')
| -rw-r--r-- | sci-libs/miopen/Manifest | 7 | ||||
| -rw-r--r-- | sci-libs/miopen/files/miopen-5.1.3-avoid-metadata-error-for-vanilla-clang.patch | 183 | ||||
| -rw-r--r-- | sci-libs/miopen/files/miopen-5.1.3-deprecate-clang-ocl.patch | 55 | ||||
| -rw-r--r-- | sci-libs/miopen/files/miopen-5.1.3-gfx1031.patch | 397 | ||||
| -rw-r--r-- | sci-libs/miopen/files/miopen-5.1.3-include-array.patch | 12 | ||||
| -rw-r--r-- | sci-libs/miopen/files/miopen-5.1.3-no-strip.patch | 17 | ||||
| -rw-r--r-- | sci-libs/miopen/miopen-5.1.3.ebuild | 109 |
7 files changed, 780 insertions, 0 deletions
diff --git a/sci-libs/miopen/Manifest b/sci-libs/miopen/Manifest index 90002c4b7555..bf01b50818af 100644 --- a/sci-libs/miopen/Manifest +++ b/sci-libs/miopen/Manifest @@ -7,8 +7,15 @@ AUX miopen-4.3.0-strip-xnack-in-flags.patch 1032 BLAKE2B 6cd0d65676edd394e8b8e06 AUX miopen-5.0.2-gfx1031.patch 11488 BLAKE2B 25b3136015309e8a06de6fcac7864f90d146beb50674db9810ed568b1ff1e8a761bf3a2a96f74ca4e215c1c670df0fcf09bb60c51d54122a063c8c8ee4059ebf SHA512 ebc1fb32aeef147366440e4ded002a833cc2f7b621b5cb2059ec5b3cc5786167c3f5955d04b0e6e3d9cfab86fe6f414e2e74a15db0090429fd97be3671c7445c AUX miopen-5.0.2-no-strip.patch 670 BLAKE2B 20a9c14b5df7d1d964ade240fc5c77ff3d119bded10f3b60e6bb17b98db116d3b7e954a6fb9ff6b5bf47243ca0e2c6160e63f27113a0663ee62eeac438c50f8d SHA512 342238b750bdf8210ccb6fc2e4839ce8d236efeec8bbf23b8fb70b0e31027aac39a8f443f0a38835d59365bc9d55bb0ecb210695c9c6488cd8e765c42754b674 AUX miopen-5.0.2-strip-xnack-in-flags.patch 1205 BLAKE2B d1ba065387050b51288dba37631adea42398b2f7cd8b604e5c305d3491fb661dce8c448e31db42c70f8b6e4dec84c4214ce869dd1419eb67c2ea4c33f18849b5 SHA512 ad32335831c7d1057b53469aa2d1f8f3366b23ccdc71cb7e9a4e7813e7da7e9bb73e62465f1086391653538ccbab2f662baf2446e10e8dbc0f0db631d7368a4d +AUX miopen-5.1.3-avoid-metadata-error-for-vanilla-clang.patch 10748 BLAKE2B 582b362a54d9ee00c8f6062b4ec16bebdaea7734adadcd5f6201801cb8ce3a7ca64e8148a10f2f3ab682560ced1f037b217bdcd161869d80a8489f4d56bd5241 SHA512 08cd489d23a48d0b227dacc514e3f84999c1e9a0ac3dd5424d9376c77c4d745c65f48c0f470cb5e65c20e9dbbfcb457995301a74bfe27a81604bf89213a2f1d2 +AUX miopen-5.1.3-deprecate-clang-ocl.patch 2884 BLAKE2B 4f45f66d87dea51d9f87878f52520863b36defa1403055a580e5f33c18e3de08daad1b3d3d5fad06d7d3f16da4f0945390872f573f168d1893f5f635070f1463 SHA512 c4835559a4005ffd2aac7afb38213d06857ab1a8d4bc5bd5c150b82226671fb1d66c6ef1d4eab2106ad11c475007af441182ca26f659a3fd779cd7d4f5c90247 +AUX miopen-5.1.3-gfx1031.patch 22036 BLAKE2B b3072b4f21162324c4ac63a020e5c448ed2df02032a7bf4e2db6f2dfc9742fddae2a7a780cf197f6c4e08123a3fa19b4f76ac22ae33d89f53adec33151eaa6b7 SHA512 8ca0c6f2aa2ffda460cf895be8d9a68aff581f6c388ba5145aa40a2b952f896901c74a24c9a03aa1c693f5f798909922bc8f056efe746f22feba8adccd0dd3d1 +AUX miopen-5.1.3-include-array.patch 343 BLAKE2B 993a8f3e301f4c16088ca216b2d820b17ceb42eeda4550b4291a70e00592281f992bf184407d024a0f8b2207a26fef2d56c26ef95e7a452d3b97090e1335a404 SHA512 887ac13f25c1bfd524bb2a3be701bf688f002a95ef8218cb798e41feaeb0678dbbd2ffaf905c59070a125001093bf5ede1bb6a77fdfe6f7a4ff19cbccfdc6624 +AUX miopen-5.1.3-no-strip.patch 608 BLAKE2B 3173d95219542a48c6dcd2b87e7157dcf483ff714c3d1c1c47aab8fc8af7b1fd4f6ac52b9d8e892cc25bff973597f4a05426fa4714b0764733b810867af59cfe SHA512 201357718203ee9a4062362541c483e069321b825708c6aa0e8892cc7722bcd520808797307e320df34c4fbf0003198f7a46717d5ff3819b1e2c3a5a3a045725 DIST MIOpen-4.3.0.tar.gz 59405686 BLAKE2B fe91bd91a74023866883d6d0a2a8071a8fe40c4cff2fb4ef58fc6e343a05ac2a731f73e657f4d183ade4e5b7c1fbbe41f3f032918f6e50cb713073aee8d97dc5 SHA512 a8615b7738acfbc7f68d9417b0746c62630d2b48fb8485fafba4db65a4b277a8230f601d249d7e54f89ba25c14176429ca76ad8365a437b09d836b031b0c4fbb DIST MIOpen-5.0.2.tar.gz 76294827 BLAKE2B 7b2a1f0e675793aee4a0fa2a270caac8332cda36c8f04cee483cc2882ed987b6e676e9c24a1acf4976a16a10f922b1a6263470b419aa88a29cfcb6d6d4b4cc29 SHA512 a581b45220797904db3e4dd3840f2ef96085f00baf8187c5ab574325a66da4f599dee6496457bb1cc32825b57a13fb0ef35a2ef1bd2a5f449c7e7b9fa64b27d1 +DIST MIOpen-5.1.3.tar.gz 88118329 BLAKE2B d24722ffc5f5dab6d6a1de2ce34193ad2f25c9a2562e38c52e010a29870f01d9ea1c56970ba0601a088c8286e97958ee95d0da27fc8082126dd2ebe5ccb36b70 SHA512 a14e28cfcb12e5061e0e7b999ef3e67fa0a0e897e31bc50e7288b8a23eb1791312e33d3b697021c2b654ccc065ae1b046c1cfd77ba8e04b0f3e87e9cc0626dcd EBUILD miopen-4.3.0.ebuild 2200 BLAKE2B 2f13af068c451192140c7df7957440e02a7f7721036c344748cb83b8f934d81ea43415625d771fe8cd73d436ad302e61dbfd1477933277ab43e9cf9b58af268c SHA512 28fcd69a452a824de2b0bba8e960940506b9b1f9e9c4eef39c842ee26b0aa6ed6e8349cb55a83658ab3fb7f4cb763092857e02216e86da3aaff54b15826a7807 EBUILD miopen-5.0.2.ebuild 2221 BLAKE2B affa8b448a9b67b91921c1fd1390ecece736fb12e40bf24f614b8460d88134810dae2136fc74a5eae1adcf41f8bfd761354679871d834fd340051745ba449ad4 SHA512 3b3197f2b73ba51e8bf095f35d7e4c11e7baefb9de74c67153c99ab1ce0c589f425a3ee40e1779e1c35eaffdbbb69b02e7ee6af0757684f8e4478647015b5250 +EBUILD miopen-5.1.3.ebuild 3130 BLAKE2B aea79cb60e5eb33d5036037e82dc7f28dd7bc3f91e8441742b570e1b915a5299df715d1416f7fc1f6db0487d31a1971f8132cda54c01c2798f0cf8c2e3cd6ab2 SHA512 6fee1d5acb2c66150c397ba5d2e536c22de1f00a32ba189b63edc1145a7a50165eb95bbd8fd9ba6ee19168473355dd3ac91595d6f3212785795a023d0c38f5d0 MISC metadata.xml 453 BLAKE2B 4c1e76cbc277e93a92a95cda405973cc72863882ab9a6729a07976e7815119e985688387c14fc5017fd6e85aaa64eaa72b900375c7f7428eb6bbb2569c5ecaa6 SHA512 0e42066a5ac8720d1b561bcf23269013d9cceff52a9ddd6a6e0af58c7d2628c335bb290e0dc6c262a0542e286f54fd07cb487dd908343fac5cfca07410df6e56 diff --git a/sci-libs/miopen/files/miopen-5.1.3-avoid-metadata-error-for-vanilla-clang.patch b/sci-libs/miopen/files/miopen-5.1.3-avoid-metadata-error-for-vanilla-clang.patch new file mode 100644 index 000000000000..3dca20f1fb12 --- /dev/null +++ b/sci-libs/miopen/files/miopen-5.1.3-avoid-metadata-error-for-vanilla-clang.patch @@ -0,0 +1,183 @@ +Together with find-sed command in 5.1.3 ebuild, this fixes +https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1731 +index 71b2cabf1..60e7fab6e 100644 +--- a/src/kernels/Conv_Winograd_v13_3_12_epilogue.inc ++++ b/src/kernels/Conv_Winograd_v13_3_12_epilogue.inc +@@ -76,7 +76,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/Conv_Winograd_v16_5_0_epilogue.inc b/src/kernels/Conv_Winograd_v16_5_0_epilogue.inc +index 36d47c862..f8f677ec6 100644 +--- a/src/kernels/Conv_Winograd_v16_5_0_epilogue.inc ++++ b/src/kernels/Conv_Winograd_v16_5_0_epilogue.inc +@@ -76,7 +76,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/Conv_Winograd_v21_1_3_metadata.inc b/src/kernels/Conv_Winograd_v21_1_3_metadata.inc +index deff81e84..ed47abea7 100644 +--- a/src/kernels/Conv_Winograd_v21_1_3_metadata.inc ++++ b/src/kernels/Conv_Winograd_v21_1_3_metadata.inc +@@ -51,7 +51,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv1x1u.s b/src/kernels/conv1x1u.s +index 5dc213546..c890d45a4 100644 +--- a/src/kernels/conv1x1u.s ++++ b/src/kernels/conv1x1u.s +@@ -1076,7 +1076,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv1x1u_bias_activ.s b/src/kernels/conv1x1u_bias_activ.s +index 1675e819a..6bbdd9936 100644 +--- a/src/kernels/conv1x1u_bias_activ.s ++++ b/src/kernels/conv1x1u_bias_activ.s +@@ -1230,7 +1230,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv1x1u_stride2.s b/src/kernels/conv1x1u_stride2.s +index c5ea1e90c..6241edcf7 100644 +--- a/src/kernels/conv1x1u_stride2.s ++++ b/src/kernels/conv1x1u_stride2.s +@@ -1162,7 +1162,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv1x1wrw.s b/src/kernels/conv1x1wrw.s +index b13b6ffa4..eb63f17c6 100644 +--- a/src/kernels/conv1x1wrw.s ++++ b/src/kernels/conv1x1wrw.s +@@ -1243,7 +1243,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv3x3wrw.s b/src/kernels/conv3x3wrw.s +index a3f73aeae..b6fb1632c 100755 +--- a/src/kernels/conv3x3wrw.s ++++ b/src/kernels/conv3x3wrw.s +@@ -1033,7 +1033,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv_3x3_wheel_alpha_v3_0b_epilogue.inc b/src/kernels/conv_3x3_wheel_alpha_v3_0b_epilogue.inc +index 358772e63..b27ad5284 100644 +--- a/src/kernels/conv_3x3_wheel_alpha_v3_0b_epilogue.inc ++++ b/src/kernels/conv_3x3_wheel_alpha_v3_0b_epilogue.inc +@@ -76,7 +76,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv_3x3_wheel_alpha_v7_0_3b_epilogue.inc b/src/kernels/conv_3x3_wheel_alpha_v7_0_3b_epilogue.inc +index d3296969a..f873ce153 100644 +--- a/src/kernels/conv_3x3_wheel_alpha_v7_0_3b_epilogue.inc ++++ b/src/kernels/conv_3x3_wheel_alpha_v7_0_3b_epilogue.inc +@@ -76,7 +76,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv_3x3_wheel_alpha_v9_0_15_epilogue.inc b/src/kernels/conv_3x3_wheel_alpha_v9_0_15_epilogue.inc +index a253cc0f9..1582d002c 100644 +--- a/src/kernels/conv_3x3_wheel_alpha_v9_0_15_epilogue.inc ++++ b/src/kernels/conv_3x3_wheel_alpha_v9_0_15_epilogue.inc +@@ -76,7 +76,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/conv_3x3_wheel_alpha_v9_2_7_epilogue.inc b/src/kernels/conv_3x3_wheel_alpha_v9_2_7_epilogue.inc +index e40ac1f60..78495e024 100644 +--- a/src/kernels/conv_3x3_wheel_alpha_v9_2_7_epilogue.inc ++++ b/src/kernels/conv_3x3_wheel_alpha_v9_2_7_epilogue.inc +@@ -76,7 +76,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/xform_bidirect_winograd_code.inc b/src/kernels/xform_bidirect_winograd_code.inc +index 724845f49..d03314ef5 100644 +--- a/src/kernels/xform_bidirect_winograd_code.inc ++++ b/src/kernels/xform_bidirect_winograd_code.inc +@@ -1566,7 +1566,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } +diff --git a/src/kernels/xform_metadata.inc b/src/kernels/xform_metadata.inc +index 960a9a2d8..83b736bee 100644 +--- a/src/kernels/xform_metadata.inc ++++ b/src/kernels/xform_metadata.inc +@@ -74,7 +74,7 @@ amdhsa.kernels: + .max_flat_workgroup_size: \wg_x + .wavefront_size: 64 + .args: +- - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: N } ++ - { .size: 4, .offset: 0, .value_kind: by_value, .value_type: i32, .name: X } + - { .size: 4, .offset: 4, .value_kind: by_value, .value_type: i32, .name: C } + - { .size: 4, .offset: 8, .value_kind: by_value, .value_type: i32, .name: H } + - { .size: 4, .offset: 12, .value_kind: by_value, .value_type: i32, .name: W } diff --git a/sci-libs/miopen/files/miopen-5.1.3-deprecate-clang-ocl.patch b/sci-libs/miopen/files/miopen-5.1.3-deprecate-clang-ocl.patch new file mode 100644 index 000000000000..93c66fd3ef05 --- /dev/null +++ b/sci-libs/miopen/files/miopen-5.1.3-deprecate-clang-ocl.patch @@ -0,0 +1,55 @@ +This is a cherry picked PR on 5.1.3, which replace clang-ocl with clang +From 98f001dfe61208af04ecf7690023efd772ee7d43 Mon Sep 17 00:00:00 2001 +From: Jehandad Khan <jahandad@gmail.com> +Date: Tue, 19 Jul 2022 17:24:05 -0500 +Subject: [PATCH] Remove clang-ocl and replace with clang + +--- + CMakeLists.txt | 7 +------ + README.md | 1 - + src/hipoc/hipoc_program.cpp | 7 ++++++- + 3 files changed, 7 insertions(+), 8 deletions(-) + +Index: MIOpen-rocm-5.1.3/CMakeLists.txt +=================================================================== +--- MIOpen-rocm-5.1.3.orig/CMakeLists.txt ++++ MIOpen-rocm-5.1.3/CMakeLists.txt +@@ -241,7 +241,7 @@ if( MIOPEN_BACKEND STREQUAL "HIP" OR MIO + # miopentensile default off + set(MIOPEN_USE_MIOPENTENSILE OFF CACHE BOOL "") + +- find_program(HIP_OC_COMPILER clang-ocl ++ find_program(HIP_OC_COMPILER clang + PATH_SUFFIXES bin + PATHS + /opt/rocm +Index: MIOpen-rocm-5.1.3/README.md +=================================================================== +--- MIOpen-rocm-5.1.3.orig/README.md ++++ MIOpen-rocm-5.1.3/README.md +@@ -14,7 +14,6 @@ MIOpen supports two programming models - + * OpenCL - OpenCL libraries and header files + * HIP - + * HIP and HCC libraries and header files +- * [clang-ocl](https://github.com/RadeonOpenCompute/clang-ocl) -- **required** + * [MIOpenGEMM](https://github.com/ROCmSoftwarePlatform/MIOpenGEMM) to enable various functionalities including transposed and dilated convolutions. This is optional on the HIP backend. Users can enable this library using the cmake configuration flag `-DMIOPEN_USE_MIOPENGEMM=On`. + * ROCm cmake modules can be installed from [here](https://github.com/RadeonOpenCompute/rocm-cmake) + * [Half](http://half.sourceforge.net/) - IEEE 754-based half-precision floating point library +Index: MIOpen-rocm-5.1.3/src/hipoc/hipoc_program.cpp +=================================================================== +--- MIOpen-rocm-5.1.3.orig/src/hipoc/hipoc_program.cpp ++++ MIOpen-rocm-5.1.3/src/hipoc/hipoc_program.cpp +@@ -255,7 +255,12 @@ void HIPOCProgramImpl::BuildCodeObjectIn + if(miopen::IsEnabled(MIOPEN_DEBUG_OPENCL_WAVE64_NOWGP{})) + params += " -mwavefrontsize64 -mcumode"; + WriteFile(src, dir->path / filename); +- dir->Execute(HIP_OC_COMPILER, params + " " + filename + " -o " + hsaco_file.string()); ++ params += " -target amdgcn-amd-amdhsa -x cl -D__AMD__=1 -O3"; ++ params += " -cl-kernel-arg-info -cl-denorms-are-zero"; ++ params += " -cl-std=CL1.2 -mllvm -amdgpu-early-inline-all"; ++ params += " -mllvm -amdgpu-internalize-symbols "; ++ params += " " + filename + " -o " + hsaco_file.string(); ++ dir->Execute(HIP_OC_COMPILER, params); + } + if(!boost::filesystem::exists(hsaco_file)) + MIOPEN_THROW("Cant find file: " + hsaco_file.string()); diff --git a/sci-libs/miopen/files/miopen-5.1.3-gfx1031.patch b/sci-libs/miopen/files/miopen-5.1.3-gfx1031.patch new file mode 100644 index 000000000000..42041f2414ed --- /dev/null +++ b/sci-libs/miopen/files/miopen-5.1.3-gfx1031.patch @@ -0,0 +1,397 @@ +Enable gfx1031 support +====================== +--- MIOpen-rocm-5.1.3/fin/src/include/conv_fin.hpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/fin/src/include/conv_fin.hpp 2022-09-01 18:36:18.630980283 +0800 +@@ -111,6 +111,10 @@ class ConvFin : public Fin + { + assert(num_cu == 72 || num_cu == 36); + } ++ else if(arch == "gfx1031") ++ { ++ assert(num_cu == 40 || num_cu == 20); ++ } + else if(arch == "gfx90a") + { + assert(num_cu == 110); +Only in MIOpen-rocm-5.1.3: patches +diff --color -upr MIOpen-rocm-5.1.3/src/composable_kernel/composable_kernel/include/utility/config.hpp gfx1031/src/composable_kernel/composable_kernel/include/utility/config.hpp +--- MIOpen-rocm-5.1.3/src/composable_kernel/composable_kernel/include/utility/config.hpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/composable_kernel/composable_kernel/include/utility/config.hpp 2022-09-01 18:36:18.634980274 +0800 +@@ -13,7 +13,7 @@ + // GPU target + // should enable one and only one GPU target + #if !(defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) || defined(CK_AMD_GPU_GFX906) || \ +- defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) || defined(CK_AMD_GPU_GFX1030)) ++ defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) || defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031)) + #error Need to define (only) one GPU target + #endif + +@@ -29,7 +29,7 @@ + #if defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) || defined(CK_AMD_GPU_GFX906) || \ + defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) + #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000 +-#elif defined(CK_AMD_GPU_GFX1030) ++#elif (defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031)) + #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000 + #endif + +@@ -37,7 +37,7 @@ + #if defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) + #define CK_USE_AMD_V_MAC_F32 + #elif defined(CK_AMD_GPU_GFX906) || defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90a) || \ +- defined(CK_AMD_GPU_GFX1030) ++ (defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031)) + #define CK_USE_AMD_V_FMAC_F32 + #define CK_USE_AMD_V_DOT2_F32_F16 + #define CK_USE_AMD_V_DOT4_I32_I8 +diff --color -upr MIOpen-rocm-5.1.3/src/include/miopen/solver/ck_utility_common.hpp gfx1031/src/include/miopen/solver/ck_utility_common.hpp +--- MIOpen-rocm-5.1.3/src/include/miopen/solver/ck_utility_common.hpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/include/miopen/solver/ck_utility_common.hpp 2022-09-01 18:36:18.638980266 +0800 +@@ -54,6 +54,7 @@ static inline bool is_ck_supported_hardw + StartsWith(handle.GetDeviceName(), "gfx908") || + StartsWith(handle.GetDeviceName(), "gfx90a") || + StartsWith(handle.GetDeviceName(), "gfx1030"); ++ StartsWith(handle.GetDeviceName(), "gfx1031"); + } + + static inline bool is_support_amd_buffer_atomic_fadd(const std::string& device_name) +@@ -83,6 +84,8 @@ static inline auto get_ck_common_compile + compiler_flag << " -DCK_AMD_GPU_GFX90A"; + else if(StartsWith(device_name, "gfx1030")) + compiler_flag << " -DCK_AMD_GPU_GFX1030"; ++ else if(StartsWith(device_name, "gfx1031")) ++ compiler_flag << " -DCK_AMD_GPU_GFX1031"; + + // buffer atomic-fadd + compiler_flag << " -DCK_USE_AMD_BUFFER_ATOMIC_FADD=" +diff --color -upr MIOpen-rocm-5.1.3/src/include/miopen/solver/implicitgemm_util.hpp gfx1031/src/include/miopen/solver/implicitgemm_util.hpp +--- MIOpen-rocm-5.1.3/src/include/miopen/solver/implicitgemm_util.hpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/include/miopen/solver/implicitgemm_util.hpp 2022-09-01 18:36:18.638980266 +0800 +@@ -469,7 +469,7 @@ static inline bool is_use_amd_buffer_loa + { + #if WORKAROUND_MIOPEN_ISSUE_557 + const auto device_name = ctx.GetStream().GetDeviceName(); +- return !StartsWith(device_name, "gfx1030"); ++ return !StartsWith(device_name, "gfx1030") && !StartsWith(device_name, "gfx1031"); + #else + return true; + #endif +@@ -478,7 +478,7 @@ static inline bool is_use_amd_buffer_loa + static inline bool is_use_v_fmac_f32(const ConvolutionContext& ctx) + { + const auto device_name = ctx.GetStream().GetDeviceName(); +- return StartsWith(device_name, "gfx1030"); ++ return StartsWith(device_name, "gfx1030") || StartsWith(device_name, "gfx1031"); + } + + static inline bool support_amd_buffer_atomic_fadd(const std::string& device_name) +@@ -599,7 +599,8 @@ static inline bool IsComposableKernelSup + StartsWith(c.GetStream().GetDeviceName(), "gfx906") || + StartsWith(c.GetStream().GetDeviceName(), "gfx908") || + StartsWith(c.GetStream().GetDeviceName(), "gfx90a") || +- StartsWith(c.GetStream().GetDeviceName(), "gfx1030"); ++ StartsWith(c.GetStream().GetDeviceName(), "gfx1030")|| ++ StartsWith(c.GetStream().GetDeviceName(), "gfx1031"); + } + + // greatest common divisor, aka highest common factor +diff --color -upr MIOpen-rocm-5.1.3/src/kernels/batchnorm_functions.h gfx1031/src/kernels/batchnorm_functions.h +--- MIOpen-rocm-5.1.3/src/kernels/batchnorm_functions.h 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/kernels/batchnorm_functions.h 2022-09-01 18:36:18.858979772 +0800 +@@ -159,6 +159,10 @@ + #define MIO_BN_GFX1030 0 + #endif + ++#ifndef MIO_BN_GFX1031 ++#define MIO_BN_GFX1031 0 ++#endif ++ + #define UNUSED __attribute__((__unused__)) + + #if(MIO_BN_VARIANT != 4) +diff --color -upr MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormActivBwdPerAct.cl gfx1031/src/kernels/MIOpenBatchNormActivBwdPerAct.cl +--- MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormActivBwdPerAct.cl 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/kernels/MIOpenBatchNormActivBwdPerAct.cl 2022-09-01 18:36:18.858979772 +0800 +@@ -34,7 +34,7 @@ + #endif + + #define MIOPEN_USE_AMDGCN 0 +-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 ++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1 + #undef MIOPEN_USE_AMDGCN + #define MIOPEN_USE_AMDGCN 1 + #endif +diff --color -upr MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormActivBwdSpatial.cl gfx1031/src/kernels/MIOpenBatchNormActivBwdSpatial.cl +--- MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormActivBwdSpatial.cl 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/kernels/MIOpenBatchNormActivBwdSpatial.cl 2022-09-01 18:36:18.858979772 +0800 +@@ -32,7 +32,7 @@ + #endif + + #define MIOPEN_USE_AMDGCN 0 +-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 ++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1 + #undef MIOPEN_USE_AMDGCN + #define MIOPEN_USE_AMDGCN 1 + #endif +diff --color -upr MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl gfx1031/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl +--- MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl 2022-09-01 18:36:18.858979772 +0800 +@@ -33,7 +33,7 @@ + #endif + + #define MIOPEN_USE_AMDGCN 0 +-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 ++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1 + #undef MIOPEN_USE_AMDGCN + #define MIOPEN_USE_AMDGCN 1 + #endif +diff --color -upr MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormBwdSpatial.cl gfx1031/src/kernels/MIOpenBatchNormBwdSpatial.cl +--- MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormBwdSpatial.cl 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/kernels/MIOpenBatchNormBwdSpatial.cl 2022-09-01 18:36:18.858979772 +0800 +@@ -33,7 +33,7 @@ + #endif + + #define MIOPEN_USE_AMDGCN 0 +-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 ++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1 + #undef MIOPEN_USE_AMDGCN + #define MIOPEN_USE_AMDGCN 1 + #endif +diff --color -upr MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl gfx1031/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl +--- MIOpen-rocm-5.1.3/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl 2022-09-01 18:36:18.858979772 +0800 +@@ -33,7 +33,7 @@ + #endif + + #define MIOPEN_USE_AMDGCN 0 +-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 ++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1 + #undef MIOPEN_USE_AMDGCN + #define MIOPEN_USE_AMDGCN 1 + #endif +diff --color -upr MIOpen-rocm-5.1.3/src/md_graph.cpp gfx1031/src/md_graph.cpp +--- MIOpen-rocm-5.1.3/src/md_graph.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/md_graph.cpp 2022-09-01 18:36:18.630980283 +0800 +@@ -738,8 +738,8 @@ void FusionMDGraph::InitConv(FusionMDGra + + add_v21_wino("gfx9", {"gfx900", "gfx906", "gfx908", "gfx90a"}, 1); + add_v21_wino("gfx9", {"gfx900", "gfx906", "gfx908", "gfx90a"}, 2); +- add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030"}, 1); +- add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030"}, 2); ++ add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030", "gfx1031"}, 1); ++ add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030", "gfx1031"}, 2); + } + } + +diff --color -upr MIOpen-rocm-5.1.3/src/ocl/fusionopbiasbnactivocl.cpp gfx1031/src/ocl/fusionopbiasbnactivocl.cpp +--- MIOpen-rocm-5.1.3/src/ocl/fusionopbiasbnactivocl.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/ocl/fusionopbiasbnactivocl.cpp 2022-09-01 18:36:18.634980274 +0800 +@@ -392,7 +392,8 @@ miopenStatus_t BatchNormBwdTrainFusionOp + " -DMIO_BN_USESAVED=" + std::to_string(static_cast<int>(true)) + + " -DMIO_BN_VARIANT=" + std::to_string(variant) + + " -DMIO_BN_CBA_WRITE_INTERMEDIATE=" + std::to_string(0) + +- " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0"); ++ " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0") + ++ " -DMIO_BN_GFX1031=" + ((handle.GetDeviceName() == "gfx1031") ? "1" : "0"); + + compile_config += add; + MIOPEN_LOG_I2(add); +@@ -607,7 +608,8 @@ miopenStatus_t BatchNormFwdTrainFusionOp + " -DMIO_SAVE_MEAN_VARIANCE=" + (saveBatchStats ? "1" : "0") + + " -DMIO_RUNNING_RESULT=" + ((savePopStats) ? "1" : "0") + + " -DMIO_BN_VARIANT=" + std::to_string(variant) + +- " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0"); ++ " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0") + ++ " -DMIO_BN_GFX1031=" + ((handle.GetDeviceName() == "gfx1031") ? "1" : "0"); + + compile_config += add; + MIOPEN_LOG_I2(add); +diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/backward_per_activation.cpp gfx1031/src/solver/batchnorm/backward_per_activation.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/backward_per_activation.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/backward_per_activation.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -113,6 +113,7 @@ BnBwdTrainingPerActivation::GetSolution( + {"MIO_BN_GRP1", ylocalsize}, + {"MIO_BN_GRP2", zlocalsize}, + {"MIO_BN_GFX1030", ((handle.GetDeviceName() == "gfx1030") ? "1" : "0")}, ++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")}, + }; + + kernel.comp_options = build_params.GenerateFor(kbp::OpenCL{}); +diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/backward_spatial_multiple.cpp gfx1031/src/solver/batchnorm/backward_spatial_multiple.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/backward_spatial_multiple.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/backward_spatial_multiple.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -210,6 +210,7 @@ ConvSolution BnBwdTrainingSpatialMultipl + {"MIO_BN_GRP1", ylocalsize}, + {"MIO_BN_GRP2", zlocalsize}, + {"MIO_BN_GFX1030", ((handle.GetDeviceName() == "gfx1030") ? "1" : "0")}, ++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")}, + {"MIO_LAYOUT_NHWC", static_cast<int>(problem.IsLayoutNHWC())}, + }; + +diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/backward_spatial_single.cpp gfx1031/src/solver/batchnorm/backward_spatial_single.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/backward_spatial_single.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/backward_spatial_single.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -247,6 +247,7 @@ BnBwdTrainingSpatialSingle::GetSolution( + + build_params << KernelBuildParameters{ + {"MIO_BN_GFX1030", (handle.GetDeviceName() == "gfx1030") ? "1" : "0"}, ++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")}, + }; + + kernel.comp_options = build_params.GenerateFor(kbp::OpenCL{}); +diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_inference.cpp gfx1031/src/solver/batchnorm/forward_inference.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_inference.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/forward_inference.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -103,6 +103,7 @@ ConvSolution BnFwdInference::GetSolution + {"MIO_BN_GRP1", ylocalsize},
+ {"MIO_BN_GRP2", zlocalsize},
+ {"MIO_BN_GFX1030", ((handle.GetDeviceName() == "gfx1030") ? "1" : "0")},
++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")},
+ };
+
+ kernel.comp_options = build_params.GenerateFor(kbp::OpenCL{});
+diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_per_activation.cpp gfx1031/src/solver/batchnorm/forward_per_activation.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_per_activation.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/forward_per_activation.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -105,6 +105,7 @@ BnFwdTrainingPerActivation::GetSolution( + {"MIO_BN_GRP1", ylocalsize}, + {"MIO_BN_GRP2", zlocalsize}, + {"MIO_BN_GFX1030", ((handle.GetDeviceName() == "gfx1030") ? "1" : "0")}, ++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")}, + }; + + auto kernel = KernelInfo{}; +diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_spatial_multiple.cpp gfx1031/src/solver/batchnorm/forward_spatial_multiple.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_spatial_multiple.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/forward_spatial_multiple.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -177,6 +177,7 @@ ConvSolution BnFwdTrainingSpatialMultipl + {"MIO_BN_GRP1", ylocalsize}, + {"MIO_BN_GRP2", zlocalsize}, + {"MIO_BN_GFX1030", ((handle.GetDeviceName() == "gfx1030") ? "1" : "0")}, ++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")}, + {"MIO_LAYOUT_NHWC", static_cast<int>(problem.IsLayoutNHWC())}, + }; + +diff --color -upr MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_spatial_single.cpp gfx1031/src/solver/batchnorm/forward_spatial_single.cpp +--- MIOpen-rocm-5.1.3/src/solver/batchnorm/forward_spatial_single.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/solver/batchnorm/forward_spatial_single.cpp 2022-09-01 18:36:18.638980266 +0800 +@@ -211,6 +211,7 @@ BnFwdTrainingSpatialSingle::GetSolution( + {"MIO_BN_GRP1", ylocalsize}, + {"MIO_BN_GRP2", zlocalsize}, + {"MIO_BN_GFX1030", ((handle.GetDeviceName() == "gfx1030") ? "1" : "0")}, ++ {"MIO_BN_GFX1031", ((handle.GetDeviceName() == "gfx1031") ? "1" : "0")}, + {"MIO_LAYOUT_NHWC", static_cast<int>(problem.IsLayoutNHWC())}, + }; + +diff --color -upr MIOpen-rocm-5.1.3/src/target_properties.cpp gfx1031/src/target_properties.cpp +--- MIOpen-rocm-5.1.3/src/target_properties.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/src/target_properties.cpp 2022-09-01 18:36:18.630980283 +0800 +@@ -54,6 +54,7 @@ static std::string GetDeviceNameFromMap( + {"Vega10", "gfx900"}, + {"gfx901", "gfx900"}, + {"10.3.0 Sienna_Cichlid 18", "gfx1030"}, ++ {"10.3.1 Navi_flounder 18", "gfx1031"}, + }; + + const char* const p_asciz = miopen::GetStringEnv(MIOPEN_DEBUG_ENFORCE_DEVICE{}); +diff --color -upr MIOpen-rocm-5.1.3/test/CMakeLists.txt gfx1031/test/CMakeLists.txt +--- MIOpen-rocm-5.1.3/test/CMakeLists.txt 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/test/CMakeLists.txt 2022-09-01 18:36:19.022979405 +0800 +@@ -38,6 +38,7 @@ option( MIOPEN_TEST_GFX90A "Test on gfx9 + option( MIOPEN_TEST_GFX900 "Test on Vega10 (gfx900)" OFF ) + option( MIOPEN_TEST_GFX906 "Test on Vega20 (gfx906)" OFF ) + option( MIOPEN_TEST_GFX1030 "Test on Navi21 (gfx1030)" OFF ) ++option( MIOPEN_TEST_GFX1031 "Test on Navi21 (gfx1031)" OFF ) + option( MIOPEN_TEST_GPU_XNACK_ENABLED "Test as if XNACK mode is enabled" OFF ) + option( MIOPEN_TEST_CONV Off) + option( MIOPEN_TEST_DEEPBENCH Off) +@@ -74,7 +75,7 @@ endif() + # Also we do not detect GPU when target GPU for testing is specified explicitly. + set(MIOPEN_TEST_GPU_DETECTION_FAILED FALSE) + set(MIOPEN_NO_GPU FALSE) +-if(NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906 OR MIOPEN_TEST_GFX908 OR MIOPEN_TEST_GFX90A OR MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_HIP_NOGPU)) ++if(NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906 OR MIOPEN_TEST_GFX908 OR MIOPEN_TEST_GFX90A OR MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031 OR MIOPEN_TEST_HIP_NOGPU)) + find_program(ROCMINFO + NAMES rocminfo + PATHS +@@ -96,6 +97,8 @@ if(NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TES + elseif (NOT ROCMINFO_EXIT_STATUS EQUAL 0) + message(WARNING "ROCMINFO FAILED, GPU TYPE UNKNOWN. Manually set respective MIOPEN_TEST_GFX* CMake variable to specify target GPU for testing.") + set(MIOPEN_TEST_GPU_DETECTION_FAILED TRUE) ++ elseif(ROCMINFO_OUTPUT MATCHES "gfx1031") ++ set(MIOPEN_TEST_GFX1031 ON) + elseif(ROCMINFO_OUTPUT MATCHES "gfx1030") + set(MIOPEN_TEST_GFX1030 ON) + elseif(ROCMINFO_OUTPUT MATCHES "gfx900") +@@ -125,6 +128,7 @@ message(STATUS "MIOPEN_TEST_GFX906 ${MIO + message(STATUS "MIOPEN_TEST_GFX908 ${MIOPEN_TEST_GFX908}") + message(STATUS "MIOPEN_TEST_GFX90A ${MIOPEN_TEST_GFX90A}") + message(STATUS "MIOPEN_TEST_GFX1030 ${MIOPEN_TEST_GFX1030}") ++message(STATUS "MIOPEN_TEST_GFX1031 ${MIOPEN_TEST_GFX1031}") + message(STATUS "MIOPEN_TEST_GPU_XNACK_ENABLED ${MIOPEN_TEST_GPU_XNACK_ENABLED}") + message(STATUS "MIOPEN_TEST_GPU_DETECTION_FAILED ${MIOPEN_TEST_GPU_DETECTION_FAILED}") + +@@ -167,10 +171,10 @@ endmacro() + set_var_to_condition(WORKAROUND_ISSUE_1187_DEFAULT MIOPEN_TEST_GFX90A AND MIOPEN_TEST_FLOAT) + option( WORKAROUND_ISSUE_1187 "" ${WORKAROUND_ISSUE_1187_DEFAULT}) + +-set_var_to_condition(WORKAROUND_ISSUE_1148_DEFAULT MIOPEN_TEST_GFX1030 AND MIOPEN_TEST_FLOAT) ++set_var_to_condition(WORKAROUND_ISSUE_1148_DEFAULT MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031 AND MIOPEN_TEST_FLOAT) + option( WORKAROUND_ISSUE_1148 "" ${WORKAROUND_ISSUE_1148_DEFAULT}) + +-set_var_to_condition(WORKAROUND_ISSUE_1334_DEFAULT MIOPEN_TEST_GFX1030 AND MIOPEN_TEST_FLOAT) ++set_var_to_condition(WORKAROUND_ISSUE_1334_DEFAULT MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031 AND MIOPEN_TEST_FLOAT) + option( WORKAROUND_ISSUE_1334 "" ${WORKAROUND_ISSUE_1334_DEFAULT}) + + if(NOT MIOPEN_TEST_MIOTENSILE) +@@ -216,7 +220,7 @@ if (MIOPEN_NO_GPU) + test_pooling3d test_perfdb) + endif() + +-if(MIOPEN_TEST_GFX1030) ++if(MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031) + if(WORKAROUND_ISSUE_1053 AND MIOPEN_TEST_ALL) + list(APPEND SKIP_TESTS test_lrn_test) + endif() +@@ -439,7 +443,7 @@ endfunction() + # If nothing is specified, the default value is taken. + # Default: FLOAT_ENABLED HALF_DISABLED BF16_DISABLED INT8_DISABLED + # +-# GPU types: GFX900, GFX906, GFX908, GFX90A, GFX1030 ++# GPU types: GFX900, GFX906, GFX908, GFX90A, GFX1030, GFX1031 + # The option can be enabled or disabled by using '_ENABLED' and '_DISABLED' suffix. + # If nothing is specified, the default value is taken. + # Default: GFX900_ENABLED, GFX906_ENABLED, GFX908_ENABLED, GFX90A_ENABLED, GFX1030_DISABLED +@@ -571,7 +575,7 @@ function(add_custom_test NAME) + set_tests_properties(${NAME} PROPERTIES RUN_SERIAL On) + endif() + +- if( (is_gfx900_check OR is_gfx906_check OR is_gfx908_check OR is_gfx1030_check OR is_gfx90a_check) ++ if( (is_gfx900_check OR is_gfx906_check OR is_gfx908_check OR is_gfx1030_check OR is_gfx1031_check OR is_gfx90a_check) + AND is_full_check + AND is_xnack_on_check + AND (is_miotensile_check AND is_mlir_check) +diff --color -upr MIOpen-rocm-5.1.3/test/handle_test.cpp gfx1031/test/handle_test.cpp +--- MIOpen-rocm-5.1.3/test/handle_test.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/test/handle_test.cpp 2022-09-01 18:36:19.018979413 +0800 +@@ -234,7 +234,7 @@ void test_warnings(kernel_type_t kern_ty + void test_arch_name() + { + auto&& h = get_handle(); +- auto known_arch = {"gfx908", "gfx90a", "gfx906", "gfx900", "gfx803", "gfx1030"}; ++ auto known_arch = {"gfx908", "gfx90a", "gfx906", "gfx900", "gfx803", "gfx1030", "gfx1031"}; + auto this_arch = h.GetDeviceName(); + EXPECT(std::any_of( + known_arch.begin(), known_arch.end(), [&](std::string arch) { return arch == this_arch; })); +diff --color -upr MIOpen-rocm-5.1.3/test/mdgraph.cpp gfx1031/test/mdgraph.cpp +--- MIOpen-rocm-5.1.3/test/mdgraph.cpp 2022-05-08 14:08:05.000000000 +0800 ++++ gfx1031/test/mdgraph.cpp 2022-09-01 18:36:19.022979405 +0800 +@@ -222,7 +222,7 @@ struct mdgraph_driver : test_driver + auto target = h.GetTargetProperties(); + + auto wino_supported_arch = { +- "gfx1030", "gfx1012", "gfx1011", "gfx90a", "gfx908", "gfx906", "gfx900", "gfx803"}; ++ "gfx1030", "gfx1031","gfx1012", "gfx1011", "gfx90a", "gfx908", "gfx906", "gfx900", "gfx803"}; + + auto is_wino_support = !xnack_enabled && + !miopen::IsDisabled(MIOPEN_DEBUG_GCN_ASM_KERNELS{}) && diff --git a/sci-libs/miopen/files/miopen-5.1.3-include-array.patch b/sci-libs/miopen/files/miopen-5.1.3-include-array.patch new file mode 100644 index 000000000000..fc6a36d22548 --- /dev/null +++ b/sci-libs/miopen/files/miopen-5.1.3-include-array.patch @@ -0,0 +1,12 @@ +This fixes compile error upon gcc-12 libstdc++ +index 1cfb2a72c..0f4feb406 100644 +--- a/test/sequences.cpp ++++ b/test/sequences.cpp +@@ -25,6 +25,7 @@ + *******************************************************************************/ + #include "test.hpp" + #include <miopen/sequences.hpp> ++#include <array> + + namespace miopen { + namespace seq { diff --git a/sci-libs/miopen/files/miopen-5.1.3-no-strip.patch b/sci-libs/miopen/files/miopen-5.1.3-no-strip.patch new file mode 100644 index 000000000000..0d1d429dbb85 --- /dev/null +++ b/sci-libs/miopen/files/miopen-5.1.3-no-strip.patch @@ -0,0 +1,17 @@ +Don't strip for release. Let portage handle stripping. +Index: MIOpen-rocm-5.1.3/CMakeLists.txt +=================================================================== +--- MIOpen-rocm-5.1.3.orig/CMakeLists.txt ++++ MIOpen-rocm-5.1.3/CMakeLists.txt +@@ -78,11 +78,6 @@ option( BUILD_DEV "Build for development + option(MIOPEN_ENABLE_FIN "Enable the fin driver for MIOpen" OFF) + + +-# Strip symbols for release +-if(NOT WIN32 AND NOT APPLE) +- set(CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE} -s") +- set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -s") +-endif() + + rocm_setup_version(VERSION 2.16.0) + diff --git a/sci-libs/miopen/miopen-5.1.3.ebuild b/sci-libs/miopen/miopen-5.1.3.ebuild new file mode 100644 index 000000000000..a7d8b5083f60 --- /dev/null +++ b/sci-libs/miopen/miopen-5.1.3.ebuild @@ -0,0 +1,109 @@ +# Copyright 1999-2022 Gentoo Authors +# Distributed under the terms of the GNU General Public License v2 + +EAPI=8 + +ROCM_VERSION=${PV} + +inherit cmake flag-o-matic llvm rocm + +LLVM_MAX_SLOT=14 + +DESCRIPTION="AMD's Machine Intelligence Library" +HOMEPAGE="https://github.com/ROCmSoftwarePlatform/MIOpen" +SRC_URI="https://github.com/ROCmSoftwarePlatform/MIOpen/archive/rocm-${PV}.tar.gz -> MIOpen-${PV}.tar.gz" + +LICENSE="MIT" +KEYWORDS="~amd64" +SLOT="0/$(ver_cut 1-2)" + +IUSE="debug test" +RESTRICT="!test? ( test )" + +RDEPEND=" + >=dev-util/hip-5.1.3 + >=dev-db/sqlite-3.17 + sci-libs/rocBLAS:${SLOT}[${ROCM_USEDEP}] + >=dev-libs/boost-1.72 +" + +DEPEND="${RDEPEND}" + +BDEPEND="dev-libs/half:0/1 + dev-util/rocm-cmake +" + +S="${WORKDIR}/MIOpen-rocm-${PV}" + +PATCHES=( + "${FILESDIR}/${PN}-4.2.0-disable-no-inline-boost.patch" + "${FILESDIR}/${PN}-4.2.0-gcc11-numeric_limits.patch" + "${FILESDIR}/${PN}-5.0.2-strip-xnack-in-flags.patch" + "${FILESDIR}/${PN}-4.3.0-fix-interface-include-in-HIP_COMPILER_FLAGS.patch" + "${FILESDIR}/${PN}-4.3.0-enable-test.patch" + "${FILESDIR}/${PN}-5.1.3-gfx1031.patch" + "${FILESDIR}/${PN}-5.1.3-deprecate-clang-ocl.patch" + "${FILESDIR}/${PN}-5.1.3-no-strip.patch" + "${FILESDIR}/${PN}-5.1.3-include-array.patch" + "${FILESDIR}/${PN}-5.1.3-avoid-metadata-error-for-vanilla-clang.patch" +) + +src_prepare() { + cmake_src_prepare + + sed -e "s:/opt/rocm/llvm:$(get_llvm_prefix ${LLVM_MAX_SLOT}) NO_DEFAULT_PATH:" \ + -e "s:/opt/rocm/hip:$(hipconfig -p) NO_DEFAULT_PATH:" \ + -e '/set( MIOPEN_INSTALL_DIR/s:miopen:${CMAKE_INSTALL_PREFIX}:' \ + -e '/MIOPEN_TIDY_ERRORS ALL/d' \ + -i CMakeLists.txt || die + + sed -e "/rocm_install_symlink_subdir(\${MIOPEN_INSTALL_DIR})/d" -i src/CMakeLists.txt || die + sed -e "/add_test/s:--build \${CMAKE_CURRENT_BINARY_DIR}:--build ${BUILD_DIR}:" -i test/CMakeLists.txt || die + + sed -e "s:\${AMD_DEVICE_LIBS_PREFIX}/lib:${EPREFIX}/usr/lib/amdgcn/bitcode:" -i cmake/hip-config.cmake || die + + # This plus avoid-metadata-error-for-vanilla-clang.patch fix bug mentioned + # in https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1731 + find src/kernels -name "*.s" -exec \ + sed -e "s/.name: n /.name: x /g" -e "s/.name: y /.name: z /g" \ + -e "s/.name: y,/.name: z,/g" -i {} \; || die +} + +src_configure() { + if ! use debug; then + append-cflags "-DNDEBUG" + append-cxxflags "-DNDEBUG" + CMAKE_BUILD_TYPE="Release" + else + CMAKE_BUILD_TYPE="Debug" + fi + + local mycmakeargs=( + -DCMAKE_SKIP_RPATH=ON + -DAMDGPU_TARGETS="$(get_amdgpu_flags)" + -DCMAKE_INSTALL_PREFIX="${EPREFIX}/usr" + -DMIOPEN_BACKEND=HIP + -DBoost_USE_STATIC_LIBS=OFF + -DMIOPEN_USE_MLIR=OFF + -DBUILD_TESTS=$(usex test ON OFF) + -DMIOPEN_TEST_ALL=$(usex test ON OFF) + ) + + if use test; then + for gpu_target in ${AMDGPU_TARGETS}; do + mycmakeargs+=( -DMIOPEN_TEST_${gpu_target^^}=ON ) + done + fi + + addpredict /dev/kfd + addpredict /dev/dri/ + append-cxxflags "--rocm-path=$(hipconfig -R)" + append-cxxflags "--hip-device-lib-path=${EPREFIX}/usr/lib/amdgcn/bitcode" + CXX="$(get_llvm_prefix ${LLVM_MAX_SLOT})/bin/clang++" cmake_src_configure +} + +src_test() { + check_amdgpu + export LD_LIBRARY_PATH="${BUILD_DIR}"/lib + MAKEOPTS="-j1" cmake_src_test +} |