Skip to content

Commit

Permalink
Correctly tag matmul tests requiring sm80 (iree-org#14173)
Browse files Browse the repository at this point in the history 
Also removes list comprehensions over only one element, which were
making this hard to parse.

Part of iree-org#14169
  • Loading branch information
@GMNGeoffrey @jvstokes
GMNGeoffrey authored and jvstokes committed Jun 25, 2023
1 parent 76cb59a commit 19884b7
Show file tree
Hide file tree
Showing 2 changed files with 30 additions and 40 deletions.
60 changes: 25 additions & 35 deletions tests/e2e/matmul/BUILD.bazel
View file
Original file line number Diff line number Diff line change
Expand Up @@ -218,13 +218,13 @@ X86_64_AVX512_VNNI = X86_64_AVX512_BASE + [
"large",
]]

[iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_%s" % compilation_info,
iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_LLVMGPUMatmulSimt",
generator = ":generate_e2e_matmul_tests",
generator_args = [
"--lhs_rhs_type=f32",
"--shapes=gpu_large_aligned",
"--compilation_info=%s" % compilation_info,
"--compilation_info=LLVMGPUMatmulSimt",
],
tags = [
# CUDA cuInit fails with sanitizer on.
Expand All @@ -238,38 +238,34 @@ X86_64_AVX512_VNNI = X86_64_AVX512_BASE + [
("cuda", "cuda"),
],
trace_runner = "//tools:iree-e2e-matmul-test",
) for compilation_info in [
"LLVMGPUMatmulSimt",
]]
)

# Testing Ampere + TensorCore path.
# WMMA TensorCore(F32): wmma.161616.f32.tf32
[iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_%s" % compilation_info,
iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_LLVMGPUMatmulTensorCore",
compiler_flags = [
"--iree-hal-cuda-llvm-target-arch=sm_80",
],
generator = ":generate_e2e_matmul_tests",
generator_args = [
"--lhs_rhs_type=f32",
"--shapes=gpu_large_aligned",
"--compilation_info=%s" % compilation_info,
"--compilation_info=LLVMGPUMatmulTensorCore",
],
tags = [
# CUDA cuInit fails with sanitizer on.
"noasan",
"nomsan",
"notsan",
"noubsan",
"requires-gpu-nvidia",
"requires-gpu-sm80",
],
target_backends_and_drivers = [
("cuda", "cuda"),
],
trace_runner = "//tools:iree-e2e-matmul-test",
) for compilation_info in [
"LLVMGPUMatmulTensorCore",
]]
)

iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_unaligned",
Expand All @@ -287,7 +283,7 @@ iree_generated_trace_runner_test(
"nomsan",
"notsan",
"noubsan",
"requires-gpu-nvidia",
"requires-gpu-sm80",
],
target_backends_and_drivers = [
("cuda", "cuda"),
Expand All @@ -296,88 +292,82 @@ iree_generated_trace_runner_test(
)

# MMA.SYNC TensorCore(F32): mma.sync.1688.f32.t32
[iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_mma_sync_%s" % compilation_info,
iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f32_gpu_large_mma_sync_LLVMGPUMatmulTensorCoreMmaSync",
compiler_flags = [
"--iree-hal-cuda-llvm-target-arch=sm_80",
],
generator = ":generate_e2e_matmul_tests",
generator_args = [
"--lhs_rhs_type=f32",
"--shapes=gpu_large_aligned",
"--compilation_info=%s" % compilation_info,
"--compilation_info=LLVMGPUMatmulTensorCoreMmaSync",
],
tags = [
# CUDA cuInit fails with sanitizer on.
"noasan",
"nomsan",
"notsan",
"noubsan",
"requires-gpu-nvidia",
"requires-gpu-sm80",
],
target_backends_and_drivers = [
("cuda", "cuda"),
],
trace_runner = "//tools:iree-e2e-matmul-test",
) for compilation_info in [
"LLVMGPUMatmulTensorCoreMmaSync",
]]
)

# WMMA TensorCore(F16): wmma.161616.f16.f16
[iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f16_gpu_large_%s" % compilation_info,
iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f16_gpu_large_LLVMGPUMatmulTensorCore",
compiler_flags = [
"--iree-hal-cuda-llvm-target-arch=sm_80",
],
generator = ":generate_e2e_matmul_tests",
generator_args = [
"--lhs_rhs_type=f16",
"--shapes=gpu_large_aligned",
"--compilation_info=%s" % compilation_info,
"--compilation_info=LLVMGPUMatmulTensorCore",
],
tags = [
# CUDA cuInit fails with sanitizer on.
"noasan",
"nomsan",
"notsan",
"noubsan",
"requires-gpu-nvidia",
"requires-gpu-sm80",
],
target_backends_and_drivers = [
("cuda", "cuda"),
],
trace_runner = "//tools:iree-e2e-matmul-test",
) for compilation_info in [
"LLVMGPUMatmulTensorCore",
]]
)

# MMA.SYNC TensorCore(F16): mma.sync.161616.f16.f16
[iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f16_gpu_large_mma_sync_%s" % compilation_info,
iree_generated_trace_runner_test(
name = "e2e_matmul_direct_f16_gpu_large_mma_sync_LLVMGPUMatmulTensorCoreMmaSync",
compiler_flags = [
"--iree-hal-cuda-llvm-target-arch=sm_80",
],
generator = ":generate_e2e_matmul_tests",
generator_args = [
"--lhs_rhs_type=f16",
"--shapes=gpu_large_aligned",
"--compilation_info=%s" % compilation_info,
"--compilation_info=LLVMGPUMatmulTensorCoreMmaSync",
],
tags = [
# CUDA cuInit fails with sanitizer on.
"noasan",
"nomsan",
"notsan",
"noubsan",
"requires-gpu-nvidia",
"requires-gpu-sm80",
],
target_backends_and_drivers = [
("cuda", "cuda"),
],
trace_runner = "//tools:iree-e2e-matmul-test",
) for compilation_info in [
"LLVMGPUMatmulTensorCoreMmaSync",
]]
)

[iree_generated_trace_runner_test(
name = "e2e_matmul_direct_%s_large_split_k" % lhs_rhs_type,
Expand Down
10 changes: 5 additions & 5 deletions tests/e2e/matmul/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -407,7 +407,7 @@ iree_generated_trace_runner_test(
"nomsan"
"notsan"
"noubsan"
"requires-gpu-nvidia"
"requires-gpu-sm80"
)

iree_generated_trace_runner_test(
Expand All @@ -431,7 +431,7 @@ iree_generated_trace_runner_test(
"nomsan"
"notsan"
"noubsan"
"requires-gpu-nvidia"
"requires-gpu-sm80"
)

iree_generated_trace_runner_test(
Expand All @@ -456,7 +456,7 @@ iree_generated_trace_runner_test(
"nomsan"
"notsan"
"noubsan"
"requires-gpu-nvidia"
"requires-gpu-sm80"
)

iree_generated_trace_runner_test(
Expand All @@ -481,7 +481,7 @@ iree_generated_trace_runner_test(
"nomsan"
"notsan"
"noubsan"
"requires-gpu-nvidia"
"requires-gpu-sm80"
)

iree_generated_trace_runner_test(
Expand All @@ -506,7 +506,7 @@ iree_generated_trace_runner_test(
"nomsan"
"notsan"
"noubsan"
"requires-gpu-nvidia"
"requires-gpu-sm80"
)

iree_generated_trace_runner_test(
Expand Down

0 comments on commit 19884b7

Please sign in to comment.