Have fp8 GEMMs go to hipblaslt only#3990
Merged
Conversation
… into fp8_hipblaslt_only
Contributor
spolifroni-amd
left a comment
spolifroni-amd
left a comment
There was a problem hiding this comment.
Nothing to review for docs
ahsan-ca
reviewed
May 8, 2025
| if(string_value_of(MIGRAPHX_SET_GEMM_PROVIDER{}) == "rocblas" or gpu::gfx_default_rocblas()) | ||
|
|
||
| // disable dot & quant_dot if no hipblaslt | ||
| if(not hipblaslt_supported()) |
Contributor
There was a problem hiding this comment.
Would we still not want to check for gpu::gfx_default_rocblas() for the case when hipblaslt is supported for the arch but we default to use rocblas for it?
Collaborator
Author
There was a problem hiding this comment.
With next rocm release rocblas is removing all fp8 support so only need to check we if have hipblaslt for fp8.
Contributor
There was a problem hiding this comment.
I was wondering for the case of gfx90a for which hipblaslt is supported but we default to use rocblas. So hipblaslt_supported() will return true, but we may be using rocblas for it (if default has not been overridden).
Collaborator
Author
There was a problem hiding this comment.
Lowering has been changed in this PR to move all FP8 gemm instructions to hipblaslt
Contributor
There was a problem hiding this comment.
I was wondering for the case of gfx90a for which hipblaslt is supported but we default to use rocblas. So hipblaslt_supported() will return true, but we may be using rocblas for it (if default has not been overridden).
For posterity: For gfx90a even though it would default to rocblas, for fp8 it would still use hipblaslt. That is intentional.
ahsan-ca
approved these changes
May 9, 2025
Contributor
ahsan-ca
left a comment
ahsan-ca
left a comment
There was a problem hiding this comment.
LGTM. Thanks for making the change.
lakhinderwalia
approved these changes
May 14, 2025
Collaborator
This build is not recommended to merge 🔴 |
Collaborator
❌bert-mrpc-tf: ERROR - check error output2025-05-14 14:32:50.791326: I tensorflow/core/platform/cpu_feature_guard.cc:210] This TensorFlow binary is optimized to use available CPU instructions in performance-critical operations.To enable the following instructions: SSE3 SSE4.1 SSE4.2 AVX AVX2 FMA, in other operations, rebuild TensorFlow with the appropriate compiler flags. WARNING: All log messages before absl::InitializeLog() is called are written to STDERR I0000 00:00:1747251176.263869 163426 gpu_device.cc:2022] Created device /job:localhost/replica:0/task:0/device:GPU:0 with 62973 MB memory: -> device: 0, name: AMD Instinct MI250X/MI250, pci bus id: 0000:b3:00.0 WARNING: All log messages before absl::InitializeLog() is called are written to STDERR I0000 00:00:1747251177.099417 163426 mlir_graph_optimization_pass.cc:401] MLIR V1 optimization pass is not enabled 2025-05-14 14:33:06.724724: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.724896: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.724943: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.724991: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.725042: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.725091: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.725133: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc 2025-05-14 14:33:06.725181: E external/local_xla/xla/service/gpu/llvm_gpu_backend/gpu_backend_lib.cc:250] bitcode module is required by this HLO module but was not found at ./opencl.bc error: Failure when generating HSACO error: Failure when generating HSACO error: Failure when generating HSACO error: Failure when generating HSACO error: Failure when generating HSACO error: Failure when generating HSACO error: Failure when generating HSACO error: Failure when generating HSACO 2025-05-14 14:33:06.726065: E tensorflow/compiler/mlir/tools/kernel_gen/tf_framework_c_interface.cc:228] INTERNAL: Generating device code failed. 2025-05-14 14:33:06.727112: W tensorflow/core/framework/op_kernel.cc:1829] UNKNOWN: JIT compilation failed. 2025-05-14 14:33:06.727130: I tensorflow/core/framework/local_rendezvous.cc:405] Local rendezvous is aborting with status: UNKNOWN: JIT compilation failed. [[{{node import/bert/embeddings/LayerNorm/moments/SquaredDifference}}]] 2025-05-14 14:33:06.727141: I tensorflow/core/framework/local_rendezvous.cc:405] Local rendezvous is aborting with status: UNKNOWN: JIT compilation failed. [[{{node import/bert/embeddings/LayerNorm/moments/SquaredDifference}}]] [[import/loss/output/_21]] 2025-05-14 14:33:06.727178: I tensorflow/core/framework/local_rendezvous.cc:424] Local rendezvous recv item cancelled. Key hash: 11217777527359497193 Traceback (most recent call last): File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 1407, in _do_call return fn(*args) File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 1390, in _run_fn return self._call_tf_sessionrun(options, feed_dict, fetch_list, File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 1483, in _call_tf_sessionrun return tf_session.TF_SessionRun_wrapper(self._session, options, feed_dict, tensorflow.python.framework.errors_impl.UnknownError: 2 root error(s) found. (0) UNKNOWN: JIT compilation failed. [[{{node import/bert/embeddings/LayerNorm/moments/SquaredDifference}}]] [[import/loss/output/_21]] (1) UNKNOWN: JIT compilation failed. [[{{node import/bert/embeddings/LayerNorm/moments/SquaredDifference}}]] 0 successful operations. 0 derived errors ignored. During handling of the above exception, another exception occurred: Traceback (most recent call last): File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 340, in main() File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 324, in main y_out = sess.run(y, feed_dict=tf_dict) File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 977, in run result = self._run(None, fetches, feed_dict, options_ptr, File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 1220, in _run results = self._do_run(handle, final_targets, final_fetches, File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 1400, in _do_run return self._do_call(_run_fn, feeds, fetches, targets, options, File "/usr/local/lib/python3.10/dist-packages/tensorflow/python/client/session.py", line 1426, in _do_call raise type(e)(node_def, op, message) # pylint: disable=no-value-for-parameter tensorflow.python.framework.errors_impl.UnknownError: Graph execution error: Detected at node 'import/bert/embeddings/LayerNorm/moments/SquaredDifference' defined at (most recent call last): Node: 'import/bert/embeddings/LayerNorm/moments/SquaredDifference' Detected at node 'import/bert/embeddings/LayerNorm/moments/SquaredDifference' defined at (most recent call last): Node: 'import/bert/embeddings/LayerNorm/moments/SquaredDifference' 2 root error(s) found. (0) UNKNOWN: JIT compilation failed. [[{{node import/bert/embeddings/LayerNorm/moments/SquaredDifference}}]] [[import/loss/output/_21]] (1) UNKNOWN: JIT compilation failed. [[{{node import/bert/embeddings/LayerNorm/moments/SquaredDifference}}]] 0 successful operations. 0 derived errors ignored. Original stack trace for 'import/bert/embeddings/LayerNorm/moments/SquaredDifference': 🔴unet: FAILED: MIGraphX is not within tolerance - check verbose output🔴bert_large_uncased_fp16: FAILED: MIGraphX is not within tolerance - check verbose output❌llama2_7b: ERROR - check error outputTraceback (most recent call last):File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 340, in main() File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 205, in main model = migraphx.parse_onnx(model_name, default_dim_value=batch) RuntimeError: /src/AMDMIGraphX/src/onnx/onnx_parser.cpp:265: parse_from: PARSE_FROM: Failed reading onnx file: /new-saved-models/llama2_7b/decoder_model.onnx ❌qwen1.5-7b: ERROR - check error outputusage: accuracy_checker.py [-h] [--onnx ONNX] [--tf TF] [--provider PROVIDER][--batch BATCH] [--fill1] [--fill0] [--fp16] [--argmax] [--verbose] [--tolerance TOLERANCE] [--input-dim INPUT_DIM] [--target TARGET] [--ort-run] [--ort-logging] [--disable-offload-copy] [--disable-fast-math] [--exhaustive_tune] accuracy_checker.py: error: unrecognized arguments: input_ids attention_mask position_ids 1 256 @attention_mask 1 256 @position_ids 1 256 ❌phi3-3.8b: ERROR - check error outputusage: accuracy_checker.py [-h] [--onnx ONNX] [--tf TF] [--provider PROVIDER][--batch BATCH] [--fill1] [--fill0] [--fp16] [--argmax] [--verbose] [--tolerance TOLERANCE] [--input-dim INPUT_DIM] [--target TARGET] [--ort-run] [--ort-logging] [--disable-offload-copy] [--disable-fast-math] [--exhaustive_tune] accuracy_checker.py: error: unrecognized arguments: input_ids attention_mask position_ids 1 256 @attention_mask 1 256 @position_ids 1 256 ❌mask-rcnn: ERROR - check error outputusage: accuracy_checker.py [-h] [--onnx ONNX] [--tf TF] [--provider PROVIDER][--batch BATCH] [--fill1] [--fill0] [--fp16] [--argmax] [--verbose] [--tolerance TOLERANCE] [--input-dim INPUT_DIM] [--target TARGET] [--ort-run] [--ort-logging] [--disable-offload-copy] [--disable-fast-math] [--exhaustive_tune] accuracy_checker.py: error: unrecognized arguments: 3 800 800 ❌#whisper-large-encoder: ERROR - check error outputTraceback (most recent call last):File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 340, in main() File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 205, in main model = migraphx.parse_onnx(model_name, default_dim_value=batch) RuntimeError: /src/AMDMIGraphX/src/include/migraphx/op/convolution.hpp:100: normalize_compute_shape: CONVOLUTION: mismatched channel numbers |
causten
pushed a commit
that referenced
this pull request
May 26, 2025
Will always use hipblaslt for fp8 GEMMs with this PR. FP8 will be eliminated if hipblaslt is not avaliable.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
7 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.