Merge branch 'develop' into ubuntu-22.04-default

.github/workflows/ci.yaml

@@ -55,6 +55,7 @@ jobs:
           -DMIGRAPHX_ENABLE_GPU=On \
           -DMIGRAPHX_ENABLE_CPU=On \
           -DMIGRAPHX_ENABLE_FPGA=On \
+          -DMIGRAPHX_ENABLE_MLIR=On \
           -DBUILD_DEV=On \
           -DROCM_ENABLE_GH_ANNOTATIONS=On \
           -DCLANG_TIDY_DEPEND_ON_TARGET=Off \
@@ -169,6 +170,71 @@ jobs:
         git config --global --add safe.directory /data
         python3 tools/format.py origin/${{ github.event_name == 'pull_request' && github.base_ref || 'develop' }}
 
+  sles:
+    runs-on: ROCM-Ubuntu
+    steps:
+    - uses: actions/checkout@v3
+      with:
+        fetch-depth: 0
+
+    # In this step, this action saves a list of existing images,
+    # the cache is created without them in the post run.
+    # It also restores the cache if it exists.
+    - name: Docker layer cache
+      uses: jpribyl/action-docker-layer-caching@v0.1.1
+      with:
+        key: docker-layer-caching-migraphx-sles-${{hashFiles('hip-clang.docker', '**/*requirements.txt', '**/install_prereqs.sh', 'rbuild.ini')}}
+        restore-keys:
+          docker-layer-caching-migraphx-sles-
+      # Ignore the failure of a step and avoid terminating the job.
+      continue-on-error: true
+
+    - name: Build the Docker image
+      run: docker build . --file tools/docker/sles.docker --tag migraphx-sles
+
+    - name: Restore cache files for ccache
+      uses: actions/cache/restore@v3 
+      id: ccache_restore
+      with:
+        path: ${{ github.workspace }}/ccache
+        key: ccache-sles-${{ github.ref }}
+        restore-keys: ccache-sles-
+
+    - name: Build migraphx
+      shell: bash -c "docker run -i -v=$GITHUB_WORKSPACE:/data -w /data migraphx-sles bash < {0}"
+      run: |
+        set -e
+        export CCACHE_COMPRESSLEVEL=10
+        export CCACHE_DIR=/data/ccache
+        export CCACHE_NOHASHDIR=true
+        export CCACHE_BASEDIR=/data
+        export CCACHE_MAXSIZE=1
+        mkdir build
+        cd build
+        CXX=/opt/rocm/llvm/bin/clang++ CC=/opt/rocm/llvm/bin/clang cmake \
+          -DBUILD_DEV=On \
+          -DCMAKE_CXX_COMPILER_LAUNCHER=/usr/local/bin/ccache \
+          -DCMAKE_C_COMPILER_LAUNCHER=/usr/local/bin/ccache \
+          ..
+        make -j$(nproc) tests driver
+
+    - name: Clear ccache cache before saving
+      if: ${{ steps.ccache_restore.outputs.cache-hit }}
+      shell: bash
+      env:
+        GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
+      run: |
+        set +x
+        gh extension install actions/gh-actions-cache --pin v1.0.1
+        gh actions-cache delete ${{ steps.ccache_restore.outputs.cache-matched-key }} --confirm
+
+    - name: Save cache files for ccache
+      uses: actions/cache/save@v3 
+      if: always()
+      with:
+        path: ${{ github.workspace }}/ccache
+        key: ccache-sles-${{ github.ref }}
+
   pyflakes:
     runs-on: ubuntu-20.04
 
@@ -274,11 +340,10 @@ jobs:
         # This path is specific to Ubuntu
         path: ${{ github.workspace }}/cget
         # Look to see if there is a cache hit for the corresponding requirements file
-        key: ${{ matrix.os }}-cget-4-${{ hashFiles('requirements.txt', 'dev-requirements.txt') }}
+        key: ${{ matrix.os }}-cget-4-${{ hashFiles('requirements.txt', 'dev-requirements.txt', 'rbuild.ini') }}
         restore-keys: ${{ matrix.os }}-cget-4-
 
     - name: Install dependencies
-      if: steps.deps_cache.outputs.cache-hit != 'true'
       run: rbuild prepare -d cget -s gh
 
     - name: Restore cache files for ccache

.github/workflows/sync-onnxrt-main.yaml

@@ -47,6 +47,7 @@ jobs:
             onnxruntime
             dependancies
             automated 
+            skip bot checks
           assignees: TedThemistokleous
           reviewers: TedThemistokleous causten
           draft: false

CHANGELOG.md

@@ -2,8 +2,49 @@
 
 Full documentation for MIGraphX is available at [MIGraphX Documentation](https://rocmdocs.amd.com/projects/AMDMIGraphX/en/latest/).
 
-## MIGraphX 2.5 for ROCm 5.5.0
+## MIGraphX 2.7 for ROCm 5.7.0
+### Added
+- Enabled hipRTC to not require dev packages for migraphx runtime and allow the ROCm install to be in a different directory than it was during build time
+- Add support for multi-target execution
+- Added Dynamic Batch support with C++/Python APIs
+- Add migraphx.create_argument to python API
+- Added dockerfile example for Ubuntu 22.04
+- Add TensorFlow supported ops in driver similar to exist onnx operator list
+- Add a MIGRAPHX_TRACE_MATCHES_FOR env variable to filter the matcher trace
+- Improved debugging by printing max,min,mean and stddev values for TRACE_EVAL = 2
+- use fast_math flag instead of ENV flag for GELU
+- Print message from driver if offload copy is set for compiled program
+### Optimizations
+- Optimized for ONNX Runtime 1.14.0
+- Improved compile times by only building for the GPU on the system
+- Improve performance of pointwise/reduction kernels when using NHWC layouts
+- Load specific version of the migraphx_py library
+- Annotate functions with the block size so the compiler can do a better job of optimizing 
+- Enable reshape on nonstandard shapes
+- Use half HIP APIs to compute max and min
+- Added support for broadcasted scalars to unsqueeze operator
+- Improved multiplies with dot operator
+- Handle broadcasts across dot and concat
+- Add verify namespace for better symbol resolution
+### Fixed
+- Resolved accuracy issues with FP16 resnet50
+- Update cpp generator to handle inf from  float
+- Fix assertion error during verify and make DCE work with tuples
+- Fix convert operation for NaNs
+- Fix shape typo in API test
+- Fix compile warnings for shadowing variable names
+- Add missing specialization for the `nullptr` for the hash function
+### Changed
+- Bumped version of half library to 5.6.0
+- Bumped CI to support rocm 5.6
+- Make building tests optional
+- replace np.bool with bool as per numpy request
+### Removed
+- Removed int8x4 rocBlas calls due to deprecation
+- removed std::reduce usage since not all OS' support it
+
 
+## MIGraphX 2.5 for ROCm 5.5.0
 ### Added
 - Y-Model feature to store tuning information with the optimized model
 - Added Python 3.10 bindings 
@@ -12,15 +53,11 @@ Full documentation for MIGraphX is available at [MIGraphX Documentation](https:/
 - Build support for ROCm MLIR
 - Added migraphx-driver flag to print optimizations in python (--python)
 - Added JIT implementation of the Gather and Pad operator which results in better handling of larger tensor sizes.
-
-
 ### Optimizations
 - Improved performance of Transformer based models
 - Improved performance of the Pad, Concat, Gather, and Pointwise operators
 - Improved onnx/pb file loading speed
 - Added general optimize pass which runs several passes such as simplify_reshapes/algebra and DCE in loop.
-
-
 ### Fixed
 - Improved parsing Tensorflow Protobuf files 
 - Resolved various accuracy issues with some onnx models
@@ -29,6 +66,5 @@ Full documentation for MIGraphX is available at [MIGraphX Documentation](https:/
 - Use --offload-arch instead of --cuda-gpu-arch for the HIP compiler
 - Changes inside JIT to use float accumulator for large reduce ops of half type to avoid overflow.
 - Changes inside JIT to temporarily use cosine to compute sine function.
-
 ### Changed
 - Changed version/location of 3rd party build dependencies to pick up fixes

CMakeLists.txt

@@ -51,6 +51,7 @@ project(migraphx LANGUAGES C CXX)
 include(CTest)
 
 find_package(ROCM REQUIRED)
+find_package(Threads REQUIRED)
 
 find_path(HALF_INCLUDE_DIR half.hpp PATH_SUFFIXES half)
 if (NOT HALF_INCLUDE_DIR)
@@ -227,14 +228,12 @@ rocm_enable_cppcheck(
         shadowVar
         shadowVariable
         unsafeClassDivZero
+        # Disable because of too many FPs
+        arithOperationsOnVoidPointer
         definePrefix:*test/include/test.hpp
         ctuOneDefinitionRuleViolation:*test/*
         useSmartPointer:*src/api/api.cpp
         useSmartPointer:*make_shared_array.hpp
-        constParameter:*src/targets/gpu/*.cpp
-        constParameter:*src/targets/gpu/*.hpp
-        # Suppress mlir_conv.cpp since this file will be deleted
-        *:*src/targets/gpu/mlir_conv.cpp
     FORCE
     INCONCLUSIVE
     RULE_FILE
@@ -251,10 +250,13 @@ rocm_enable_cppcheck(
         ${CMAKE_CURRENT_SOURCE_DIR}/src/targets/gpu/kernels/include
         ${CMAKE_CURRENT_SOURCE_DIR}/test/include
     DEFINE
+        MIGRAPHX_MLIR=1
         CPPCHECK=1
         __device__=
         __host__=
         __global__=
+    UNDEFINE
+        MIGRAPHX_USE_CLANG_TIDY
 )
 
 enable_testing()

Dockerfile

@@ -80,6 +80,9 @@ ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
 ADD rbuild.ini /rbuild.ini
 
+# Temporarily install a new cmake until switching to ubuntu 22.04
+RUN pip3 install cmake==3.22.1
+
 COPY ./tools/install_prereqs.sh /
 RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh
 RUN test -f /usr/local/hash || exit 1
@@ -101,6 +104,9 @@ RUN cget -p $PREFIX install facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cma
 RUN cget -p $PREFIX install ccache@v4.1 -DENABLE_TESTING=OFF
 RUN cget -p /opt/cmake install kitware/cmake@v3.26.4
 
+# Install MLIR
+ADD mlir-requirements.txt /mlir-requirements.txt
+RUN cget -p /usr/local install -f /mlir-requirements.txt
 
 COPY ./test/onnx/.onnxrt-commit /
 
@@ -116,9 +122,6 @@ RUN git clone --single-branch --branch ${ONNXRUNTIME_BRANCH} --recursive ${ONNXR
 
 ADD tools/build_and_test_onnxrt.sh /onnxruntime/build_and_test_onnxrt.sh
 
-# Use the /opt/cmake install because LLVM/MLIR need cmake >= 3.20
-RUN env PATH=/opt/cmake/bin:$PATH cget -p /usr/local install ROCmSoftwarePlatform/rocMLIR@1ad9d6df32acc6d29d58e8ed6710e36746d0a4d6 -DBUILD_FAT_LIBROCKCOMPILER=On
-
 ENV MIOPEN_FIND_DB_PATH=/tmp/miopen/find-db
 ENV MIOPEN_USER_DB_PATH=/tmp/miopen/user-db
 ENV LD_LIBRARY_PATH=$PREFIX/lib

Jenkinsfile

@@ -114,6 +114,10 @@ rocmtest clang_debug: rocmnode('cdna') { cmake_build ->
         cmake_build(flags: "-DCMAKE_BUILD_TYPE=release")
         stash includes: 'build/*.deb', name: 'migraphx-package'
     }
+// }, hidden_symbols: rocmnode('cdna') { cmake_build ->
+//     stage('Hidden symbols') {
+//         cmake_build(flags: "-DMIGRAPHX_ENABLE_PYTHON=Off -DMIGRAPHX_ENABLE_GPU=On -DMIGRAPHX_ENABLE_CPU=On -DCMAKE_CXX_VISIBILITY_PRESET=hidden -DCMAKE_C_VISIBILITY_PRESET=hidden")
+//     }
 }, all_targets_debug : rocmnode('cdna') { cmake_build ->
     stage('All targets Release') {
         cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DMIGRAPHX_ENABLE_GPU=On -DMIGRAPHX_ENABLE_CPU=On -DMIGRAPHX_ENABLE_FPGA=On") 

cmake/Embed.cmake

@@ -24,11 +24,7 @@
 find_program(EMBED_LD ld)
 find_program(EMBED_OBJCOPY objcopy)
 
-if(LINUX)
-    option(EMBED_USE_LD "Use ld to embed data files" ON)
-else()
-    option(EMBED_USE_LD "Use ld to embed data files" OFF)
-endif()
+option(EMBED_USE_LD "Use ld to embed data files" OFF)
 
 function(wrap_string)
     set(options)

cmake/ExportHeader.cmake

@@ -29,7 +29,10 @@ endif()
 include(GenerateExportHeader)
 
 function(migraphx_generate_export_header TARGET)
-    cmake_parse_arguments(PARSE "" "DIRECTORY" "" ${ARGN})
+    set(options)
+    set(oneValueArgs DIRECTORY)
+    set(multiValueArgs)
+    cmake_parse_arguments(PARSE "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
     if(PARSE_DIRECTORY)
         set(__directory ${PARSE_DIRECTORY})
     else()

cppcheck.rules

@@ -107,6 +107,24 @@
         <summary>Use make_shared or make_unique instead of new</summary>
     </message>
 </rule>
+<rule>
+    <tokenlist>raw</tokenlist>
+    <pattern><![CDATA[ [^\(,;{}:]+ \w+ && (\w+|\(|\+|\-|\*)]]></pattern>
+    <message>
+        <id>UseNamedLogicOperator</id>
+        <severity>style</severity>
+        <summary>Use 'and' instead of &&</summary>
+    </message>
+</rule>
+<rule>
+    <tokenlist>raw</tokenlist>
+    <pattern><![CDATA[ (if|while) \([^\)&]+&& ]]></pattern>
+    <message>
+        <id>UseNamedLogicOperator</id>
+        <severity>style</severity>
+        <summary>Use 'and' instead of &&</summary>
+    </message>
+</rule>
 <rule>
     <tokenlist>raw</tokenlist>
     <pattern><![CDATA[ \|\| ]]></pattern>

dev-requirements.txt

@@ -25,6 +25,6 @@ ROCmSoftwarePlatform/rocm-recipes
 facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cmake
 ccache@v4.1 -DENABLE_TESTING=OFF
 pcre,pfultz2/pcre@8.45 -H sha256:d6f7182602a775a7d500a0cedca6449af0400c6493951513046d17615ed0bf11
-danmar/cppcheck@2.9 -DHAVE_RULES=1
-RadeonOpenCompute/rocm-cmake@027404a8326da6e7e9338e0b81f9428660190724 --build
+danmar/cppcheck@bb2711c22a0be09efe7f1a8da3030876471026c8 -DHAVE_RULES=1 # 2.11
+RadeonOpenCompute/rocm-cmake@189d497ed185683154ae9766393b9a10ff21201f --build
 -f requirements.txt

docs/.sphinx/requirements.in

-rocm-docs-core==0.11.0
+rocm-docs-core>=0.20.0

docs/.sphinx/requirements.txt

 # Copyright (c) 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.
-
-#
-# This file is autogenerated by pip-compile with Python 3.8
-# by the following command:
-#
-#    pip-compile requirements.in
-#
 accessible-pygments==0.0.4
     # via pydata-sphinx-theme
 alabaster==0.7.13
@@ -46,7 +21,7 @@ charset-normalizer==3.1.0
     # via requests
 click==8.1.3
     # via sphinx-external-toc
-cryptography==40.0.2
+cryptography==41.0.3
     # via pyjwt
 deprecated==1.2.13
     # via pygithub
@@ -60,22 +35,16 @@ fastjsonschema==2.16.3
     # via rocm-docs-core
 gitdb==4.0.10
     # via gitpython
-gitpython==3.1.31
+gitpython==3.1.32
     # via rocm-docs-core
 idna==3.4
     # via requests
 imagesize==1.4.1
     # via sphinx
-importlib-metadata==6.4.1
-    # via sphinx
-importlib-resources==5.12.0
-    # via rocm-docs-core
 jinja2==3.1.2
     # via
     #   myst-parser
     #   sphinx
-linkify-it-py==1.0.3
-    # via myst-parser
 markdown-it-py==2.2.0
     # via
     #   mdit-py-plugins
@@ -86,7 +55,7 @@ mdit-py-plugins==0.3.5
     # via myst-parser
 mdurl==0.1.2
     # via markdown-it-py
-myst-parser[linkify]==1.0.0
+myst-parser==1.0.0
     # via rocm-docs-core
 packaging==23.1
     # via
@@ -109,8 +78,6 @@ pyjwt[crypto]==2.6.0
     # via pygithub
 pynacl==1.5.0
     # via pygithub
-pytz==2023.3
-    # via babel
 pyyaml==6.0
     # via
     #   myst-parser
@@ -120,7 +87,7 @@ requests==2.28.2
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core==0.11.0
+rocm-docs-core>=0.20.0
     # via -r requirements.in
 smmap==5.0.0
     # via gitdb
@@ -163,13 +130,7 @@ sphinxcontrib-serializinghtml==1.1.5
     # via sphinx
 typing-extensions==4.5.0
     # via pydata-sphinx-theme
-uc-micro-py==1.0.1
-    # via linkify-it-py
 urllib3==1.26.15
     # via requests
 wrapt==1.15.0
     # via deprecated
-zipp==3.15.0
-    # via
-    #   importlib-metadata
-    #   importlib-resources

examples/migraphx/cpp_parse_load_save/parse_load_save.cpp

@@ -114,7 +114,7 @@ int main(int argc, char** argv)
 char* getCmdOption(char** begin, char** end, const std::string& option)
 {
     char** itr = std::find(begin, end, option);
-    if(itr != end && ++itr != end)
+    if(itr != end and ++itr != end)
     {
         return *itr;
     }

examples/vision/cpp_mnist/mnist_inference.cpp

@@ -160,9 +160,9 @@ int main(int argc, char** argv)
     auto lengths = shape.lengths();
     auto num_results =
         std::accumulate(lengths.begin(), lengths.end(), 1, std::multiplies<size_t>());
-    float* results = reinterpret_cast<float*>(outputs[0].data());
-    float* max     = std::max_element(results, results + num_results);
-    int answer     = max - results;
+    float* results   = reinterpret_cast<float*>(outputs[0].data());
+    const float* max = std::max_element(results, results + num_results);
+    int answer       = max - results;
 
     std::cout << std::endl
               << "Randomly chosen digit: " << rand_digit << std::endl
@@ -192,12 +192,12 @@ void read_nth_digit(const int n, std::vector<float>& digit)
         for(int i = 0; i < HEIGHT * WIDTH; ++i)
         {
             unsigned char temp = 0;
-            file.read((char*)&temp, sizeof(temp));
+            file.read(reinterpret_cast<char*>(&temp), sizeof(temp));
             if(d == n)
             {
                 float data = temp / 255.0;
                 digit.push_back(data);
-                std::cout << SYMBOLS[(int)(data * 10) % 11];
+                std::cout << SYMBOLS[static_cast<int>(data * 10) % 11];
                 if((i + 1) % WIDTH == 0)
                     std::cout << std::endl;
             }

hip-clang.docker

@@ -57,5 +57,12 @@ ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
 ADD rbuild.ini /rbuild.ini
 
+# Temporarily install a new cmake until switching to ubuntu 22.04
+RUN pip3 install cmake==3.22.1
+
 COPY ./tools/install_prereqs.sh /
 RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh
+
+# Install MLIR
+ADD mlir-requirements.txt /mlir-requirements.txt
+RUN cget -p /usr/local install -f /mlir-requirements.txt

mlir-requirements.txt

+#####################################################################################
+# The MIT License (MIT)
+#
+# Copyright (c) 2015-2022 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.
+#####################################################################################
+ROCmSoftwarePlatform/rocMLIR@3657f509bfed86bb79d5c6e24aa237e48f09f9f3 -DBUILD_FAT_LIBROCKCOMPILER=On

rbuild.ini

@@ -6,7 +6,9 @@ deps =
     -f requirements.txt
 
 [gh]
-ignore = danmar/cppcheck
+ignore =
+    danmar/cppcheck
+    ROCmSoftwarePlatform/rocMLIR
 deps =
     -f dev-requirements.txt
     oneapi-src/oneDNN@v1.7

src/CMakeLists.txt

@@ -248,8 +248,6 @@ endif()
 target_link_libraries(migraphx PRIVATE -ldl)
 
 target_include_directories(migraphx SYSTEM PUBLIC $<BUILD_INTERFACE:${HALF_INCLUDE_DIR}>)
-
-find_package(Threads)
 target_link_libraries(migraphx PUBLIC Threads::Threads)
 
 find_package(nlohmann_json 3.8.0 REQUIRED)

src/api/api.cpp

@@ -899,7 +899,7 @@ migraphx_dynamic_dimensions_assign_to(migraphx_dynamic_dimensions_t output,
 
 extern "C" migraphx_status
 migraphx_dynamic_dimensions_create(migraphx_dynamic_dimensions_t* dynamic_dimensions,
-                                   const_migraphx_dynamic_dimension_t* ptr,
+                                   const const_migraphx_dynamic_dimension_t* ptr,
                                    size_t size)
 {
     auto api_error_result = migraphx::try_([&] {
@@ -1432,7 +1432,7 @@ extern "C" migraphx_status migraphx_instructions_assign_to(migraphx_instructions
 }
 
 extern "C" migraphx_status migraphx_instructions_create(migraphx_instructions_t* instructions,
-                                                        const_migraphx_instruction_t* ptr,
+                                                        const const_migraphx_instruction_t* ptr,
                                                         size_t size)
 {
     auto api_error_result = migraphx::try_([&] {

src/api/include/migraphx/migraphx.h

@@ -209,7 +209,7 @@ MIGRAPHX_C_EXPORT migraphx_status migraphx_dynamic_dimensions_assign_to(
 
 MIGRAPHX_C_EXPORT migraphx_status
 migraphx_dynamic_dimensions_create(migraphx_dynamic_dimensions_t* dynamic_dimensions,
-                                   const_migraphx_dynamic_dimension_t* ptr,
+                                   const const_migraphx_dynamic_dimension_t* ptr,
                                    size_t size);
 
 MIGRAPHX_C_EXPORT migraphx_status
@@ -377,7 +377,7 @@ MIGRAPHX_C_EXPORT migraphx_status migraphx_instructions_assign_to(
     migraphx_instructions_t output, const_migraphx_instructions_t input);
 
 MIGRAPHX_C_EXPORT migraphx_status migraphx_instructions_create(
-    migraphx_instructions_t* instructions, const_migraphx_instruction_t* ptr, size_t size);
+    migraphx_instructions_t* instructions, const const_migraphx_instruction_t* ptr, size_t size);
 
 MIGRAPHX_C_EXPORT migraphx_status migraphx_modules_destroy(migraphx_modules_t modules);