Add cmake logic to look for local gfx arch (#2018)

cmake/MigxLocalTarget.cmake

+# ####################################################################################
+# The MIT License (MIT)
+#
+# Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+# ####################################################################################
+# https://github.com/RadeonOpenCompute/rocm-cmake/blob/f0717bc1a0182b2ddc9194e3144972ab58a88e99/share/rocm/cmake/ROCMLocalTargets.cmake#L16C57-L16C57
+# once it is available from rocm-cmake  or hip, remove this
+function(rocm_local_targets_migx VARIABLE)
+    # rocm_agent_enumerator is only available on Linux platforms right now
+    find_program(_rocm_agent_enumerator rocm_agent_enumerator HINTS /opt/rocm/bin ENV ROCM_PATH)
+
+    if(NOT _rocm_agent_enumerator STREQUAL "_rocm_agent_enumerator-NOTFOUND")
+        execute_process(
+            COMMAND "${_rocm_agent_enumerator}"
+            RESULT_VARIABLE _found_agents
+            OUTPUT_VARIABLE _rocm_agents
+            ERROR_QUIET
+        )
+
+        if(_found_agents EQUAL 0)
+            cmake_policy(SET CMP0007 NEW)
+            string(REPLACE "\n" ";" _rocm_agents "${_rocm_agents}")
+            list(POP_BACK _rocm_agents)
+            list(REMOVE_DUPLICATES _rocm_agents)
+            unset(result)
+
+            foreach(agent IN LISTS _rocm_agents)
+                if(NOT agent STREQUAL "gfx000")
+                    list(APPEND result "${agent}")
+                endif()
+            endforeach()
+
+            if(result)
+                set(${VARIABLE} "${result}" PARENT_SCOPE)
+            endif()
+        endif()
+    endif()
+endfunction()

examples/migraphx/custom_op_hip_kernel/README.md

@@ -3,9 +3,11 @@ This is an example of a custom operator implementation using MIGraphX's C/C++ AP
 
 Kernels can be written in either HIP, MIOpen, or by using RocBLAS library. This particular example uses **HIP**.
 
- To build the example, ensure ROCm is installed at `/opt/rocm`. 
- 1.  `export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH`
- 2.  `cd $MIGRAPHX_SRC/examples/migraphx/custom_op_hip_kernel/`
- 3.  `mkdir build && cd build`
- 4.  `CXX=/opt/rocm/llvm/bin/clang++ cmake ..  && make`
- 5.  `./custom_op_hip_kernel`
\ No newline at end of file
+To build the example, ensure ROCm is installed at `/opt/rocm`. 
+
+ 1.  Before building this example, find out the gfx-arch of the machine by running `/opt/rocm/bin/rocminfo | grep -o -m 1 "gfx.*"`  pass this gfx-arch as the value of `-DGPU_TARGETS` flag in step 5. Let's assume for now it is MI200 architecture in that case gfx-arch would be `gfx90a`. 
+ 2.  `export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH`
+ 3.  `cd $MIGRAPHX_SRC/examples/migraphx/custom_op_hip_kernel/`
+ 4.  `mkdir build && cd build`
+ 5.  `CXX=/opt/rocm/llvm/bin/clang++ cmake .. -DGPU_TARGETS=gfx90a  && make`
+ 6.  `./custom_op_hip_kernel`
\ No newline at end of file

src/targets/gpu/CMakeLists.txt

@@ -21,8 +21,17 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 # ####################################################################################
-
 list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
+
+if(NOT WIN32)
+    include(MigxLocalTarget)
+    rocm_local_targets_migx(MIGXGPU_TARGETS)
+    if(MIGXGPU_TARGETS)
+        set(GPU_TARGETS ${MIGXGPU_TARGETS} CACHE STRING "GPU targets to compile for")
+        message(STATUS "Compiling for GPU targets : ${GPU_TARGETS}")
+    endif()
+endif()
+
 find_package(miopen)
 
 # rocblas