Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

New API to allow sending hipEvent to Tensile #989

Closed
wants to merge 2 commits into from

Conversation

mahmoodw
Copy link
Contributor

@mahmoodw mahmoodw commented Feb 8, 2020

Solution proposed to profile kernel time more accurately.

Please share your thoughts about adding in this API and if there is a less public way of introducing this API. Currently GemmStress uses functions from rocblas.h so this seemed necessary to me.

Copy link
Contributor

@leekillough leekillough left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We want to keep a strict separation of rocBLAS and Tensile, as shown in the rocSOLVER document I commented on.

So if we need to pass event handlers for timing or other stuff, it needs to be incorporated into tensile_host.cpp and maybe tensile_host.hpp. The rocblas_handle and gemm_ex functions should not call Tensile directly, or depend on Tensile implementation details.

tensile_host.hpp declares constructors for the PIMPL implementation class, which rocBLAS uses. tensile_host.cpp implements the class and its methods. rocBLAS only references tensile_host.hpp for its interface to Tensile. No internal details of Tensile are assumed outside of tensile_host.cpp, and no Tensile header files are #included outside of tensile_host.cpp in the new client (only in the legacy old client code).

There is also the old client code in gemm* and the new client code in tensile_host.cpp/hpp. We build the new client code by default. The old client code will be removed at some point (maybe soon -- see PR #990 ), so we should not rely on changes to its calls to Tensile for event handlers.

Event handlers can be added to rocblas_handle, and then tensile_host.cpp can use them to pass on to Tensile. There will need to be a new API for passing extra information to Tensile in the new client API, such as problem modifiers, instead of simply adding arguments to the ends of function parameter lists.

So I suggest:

  1. Adding start/stop event handlers to rocblas_handle.
  2. Add rocBLAS functions, either internal C++ methods not published, or in the public C API, to manipulate the event handlers.
  3. Change tensile_host.cpp to get the event handlers from the rocblas_handle, and pass them to Tensile using the new client API.
  4. Do not change the APIs of existing rocBLAS kernels to add extra arguments. New client is moving away from using explicit argument lists, and instead uses problems, constraints, modifiers, etc.
  5. Don’t worry about changing gemm* functions which call Tensile, since they are only used in the old client, and will be removed soon ( Remove old Tensile client code #990 ).

@mahmoodw mahmoodw closed this Feb 26, 2020
@mahmoodw mahmoodw mentioned this pull request Feb 26, 2020
mlse-lib-jenkins pushed a commit that referenced this pull request Nov 28, 2021
@amcamd amcamd deleted the gemm_ex_profile branch May 10, 2023 15:08
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

2 participants