Skip to content

tests: hip_prefill_wmma_gemm_parity skips on a card without matrix cores instead of failing #1148

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.

Sign up for GitHub

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

Closed
Vop3p wants to merge 1 commit into from
Closed

tests: hip_prefill_wmma_gemm_parity skips on a card without matrix cores instead of failing #1148

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -594,6 +594,7 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP)
target_include_directories(hip_prefill_wmma_gemm_parity PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src/prefill")
target_link_libraries(hip_prefill_wmma_gemm_parity PRIVATE strata_prefill strata_hip_runtime)
add_test(NAME hip_prefill_wmma_gemm_parity COMMAND hip_prefill_wmma_gemm_parity)
set_tests_properties(hip_prefill_wmma_gemm_parity PROPERTIES SKIP_RETURN_CODE 77) # no matrix cores: skipped
list(APPEND _strata_hip_tests hip_prefill_wmma_gemm_parity)
# S (STRATA_HCD_EXACT): the exact HC down kernel against hipBLASLt solution 1176 / 1177 (needs STRATA_HIPBLASLT_TUNING; exit 77 else)
add_executable(hip_prefill_hcd_exact_parity tests/hip/prefill_hcd_exact_parity.cpp)
Expand Down
8 changes: 8 additions & 0 deletions tests/hip/prefill_wmma_gemm_parity.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
#include <hip/hip_fp16.h>

#include "wmma_gemm.h"
#include "strata/kernels/gfx_arch.hpp"

#include <algorithm>
#include <cmath>
Expand Down Expand Up @@ -269,6 +270,13 @@ int main() {
HIP_CHECK(hipGetDeviceProperties(&properties, device));
std::printf("Device %d: %s (arch %s, warpSize %d)\n",
device, properties.name, properties.gcnArchName, properties.warpSize);
// No matrix cores (RDNA2 and older, gfx906): the kernel declines every shape, so the suite has nothing to compare
// and reports that as a skip (exit 77, SKIP_RETURN_CODE in CMake), as hip_prompt_attn_wmma does, not a failure.
if (!strata::kernels::gfx_arch_is_gfx11_wmma(properties.gcnArchName)) {
std::printf("WMMA GEMM parity SKIPPED: %s has no WMMA matrix cores (the kernel declines every shape here)\n",
properties.gcnArchName);
return 77;
}

hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
Expand Down