diff --git a/test/common.hpp b/test/common.hpp index ecdea847..416d336a 100644 --- a/test/common.hpp +++ b/test/common.hpp @@ -57,7 +57,7 @@ } #endif -#ifdef ROCWMMA_BENCHMARK_TESTS +#if ROCWMMA_BENCHMARK_TESTS #ifndef CHECK_RSMI_ERROR #define CHECK_RSMI_ERROR(expression, smiErrorFlag) \ if(auto status = (expression); status != RSMI_STATUS_SUCCESS) \ @@ -68,7 +68,7 @@ smiErrorFlag = true; \ } #endif -#endif +#endif // ROCWMMA_BENCHMARK_TESTS namespace rocwmma { diff --git a/test/dlrm/dlrm_kernel_base_impl.hpp b/test/dlrm/dlrm_kernel_base_impl.hpp index 522b3cc7..ca3c14a9 100644 --- a/test/dlrm/dlrm_kernel_base_impl.hpp +++ b/test/dlrm/dlrm_kernel_base_impl.hpp @@ -51,7 +51,7 @@ // Library includes -#ifdef ROCWMMA_VALIDATION_TESTS +#if ROCWMMA_VALIDATION_TESTS #include "reference.hpp" // Vanilla CPU kernel #endif // ROCWMMA_VALIDATION_TESTS @@ -156,11 +156,11 @@ namespace rocwmma mM = mK = mB = 0; mMPadded = mKPadded = 0; mRepeats = -#ifdef ROCWMMA_VALIDATION_TESTS +#if ROCWMMA_VALIDATION_TESTS 1; #else 5; -#endif +#endif // ROCWMMA_VALIDATION_TESTS mRunFlag = true; @@ -187,10 +187,10 @@ namespace rocwmma << "DataT, " << "Direction, " << "MatM, MatK, MatB, " -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS << "maxRelativeDiff, " << "tolerance, " -#endif +#endif // ROCWMMA_VALIDATION_TESTS << "elapsedMs, " << "Problem Size(GFlops), " << "TFlops/s, " @@ -206,9 +206,9 @@ namespace rocwmma << (passDirection == DlrmDirection_t::Forward ? "Forwards" : "Backwards") << ", " << mM << ", " << mK << ", " << mB << ", " -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS << "n/a, " -#endif +#endif // ROCWMMA_VALIDATION_TESTS << "n/a, n/a, n/a, n/a, SKIPPED" << std::endl; } else @@ -217,12 +217,12 @@ namespace rocwmma << (passDirection == DlrmDirection_t::Forward ? "Forwards" : "Backwards") << ", " << mM << ", " << mK << ", " << mB << ", " -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS << mMaxRelativeError << ", " -#endif +#endif // ROCWMMA_VALIDATION_TESTS << mElapsedTimeMs << ", " << mTotalGFlops << ", " << mMeasuredTFlopsPerSec << ", " << mEfficiency << ", " -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS << (mValidationResult ? "PASSED" : "FAILED") #else << "BENCH" @@ -275,7 +275,7 @@ namespace rocwmma { MatrixUtil::fillLaunchKernel( dataInstance->deviceInput().get(), mM, mK, mB); -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS dataInstance->copyDeviceToHostFwdInput(); #endif // ROCWMMA_VALIDATION_TESTS } @@ -286,7 +286,7 @@ namespace rocwmma dataInstance->deviceInput().get(), mM, mK, mB); MatrixUtil::fillLaunchKernel( dataInstance->deviceUpstreamGrad().get(), 1, gradSize, mB); -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS dataInstance->copyDeviceToHostBwdInput(); #endif // ROCWMMA_VALIDATION_TESTS } @@ -418,7 +418,7 @@ namespace rocwmma CHECK_HIP_ERROR(hipEventDestroy(startEvent)); CHECK_HIP_ERROR(hipEventDestroy(stopEvent)); -#if defined(ROCWMMA_VALIDATION_TESTS) +#if ROCWMMA_VALIDATION_TESTS // Run reference CPU kernel std::function cpuKernel; @@ -447,14 +447,14 @@ namespace rocwmma }; } cpuKernel(); -#endif +#endif // ROCWMMA_VALIDATION_TESTS } } template void DlrmKernelBase::validateResults() { -#ifdef ROCWMMA_VALIDATION_TESTS +#if ROCWMMA_VALIDATION_TESTS if(mRunFlag) { auto& dataInstance = DataStorage::instance(); diff --git a/test/gemm/gemm_common_test_params.hpp b/test/gemm/gemm_common_test_params.hpp index faf0a21e..aaacaf27 100644 --- a/test/gemm/gemm_common_test_params.hpp +++ b/test/gemm/gemm_common_test_params.hpp @@ -49,7 +49,7 @@ namespace rocwmma // Native int8 using TestTypesI8 = std::tuple< -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS std::tuple, #endif // ROCWMMA_EXTENDED_TESTS std::tuple>; @@ -62,7 +62,7 @@ namespace rocwmma // Non-native bfloat16_t using TestTypesBF16 = std::tuple< -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS std::tuple, std::tuple, #endif // ROCWMMA_EXTENDED_TESTS @@ -70,7 +70,7 @@ namespace rocwmma // Native f16 using TestTypesF16 = std::tuple< -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS std::tuple, std::tuple, #endif // ROCWMMA_EXTENDED_TESTS @@ -79,7 +79,7 @@ namespace rocwmma #if !ROCWMMA_TESTS_NO_HALF // Non-native hfloat16_t (i.e. __half) using TestTypesH16 = std::tuple< -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS std::tuple, std::tuple, #endif // ROCWMMA_EXTENDED_TESTS @@ -140,28 +140,28 @@ namespace rocwmma /// using TestLayoutsNN = -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS typename CombineOne, TestDataLayouts>::Result; #else std::tuple>; #endif // ROCWMMA_EXTENDED_TESTS using TestLayoutsNT = -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS typename CombineOne, TestDataLayouts>::Result; #else std::tuple>; #endif // ROCWMMA_EXTENDED_TESTS using TestLayoutsTN = -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS typename CombineOne, TestDataLayouts>::Result; #else std::tuple>; #endif // ROCWMMA_EXTENDED_TESTS using TestLayoutsTT = -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS typename CombineOne, TestDataLayouts>::Result; #else std::tuple>; @@ -177,7 +177,7 @@ namespace rocwmma // Aggregate combinations BlockK <= 32 using TestBlockSizes16x16SmallBlockK = std::tuple, I<16>, I<16>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<16>, I<32>> #endif // ROCWMMA_EXTENDED_TESTS @@ -186,7 +186,7 @@ namespace rocwmma // Aggregate combinations BlockK <= 64 using TestBlockSizes16x16MediumBlockK = std::tuple, I<16>, I<16>>, std::tuple, I<16>, I<32>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<16>, I<64>> #endif // ROCWMMA_EXTENDED_TESTS @@ -196,7 +196,7 @@ namespace rocwmma using TestBlockSizes16x16LargeBlockK = std::tuple, I<16>, I<16>>, std::tuple, I<16>, I<32>>, std::tuple, I<16>, I<64>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<16>, I<128>> #endif // ROCWMMA_EXTENDED_TESTS @@ -206,7 +206,7 @@ namespace rocwmma using TestBlockSizes16x16HugeBlockK = std::tuple, I<16>, I<16>>, std::tuple, I<16>, I<32>>, std::tuple, I<16>, I<64>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<16>, I<128>>, std::tuple, I<16>, I<256>> @@ -217,7 +217,7 @@ namespace rocwmma // Aggregate combinations BlockK <= 16 using TestBlockSizes32x32SmallBlockK = std::tuple, I<32>, I<8>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<32>, I<16>> #endif // ROCWMMA_EXTENDED_TESTS @@ -226,7 +226,7 @@ namespace rocwmma // Aggregate combinations BlockK <= 32 using TestBlockSizes32x32MediumBlockK = std::tuple, I<32>, I<8>>, std::tuple, I<32>, I<16>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<32>, I<32>> #endif // ROCWMMA_EXTENDED_TESTS @@ -236,7 +236,7 @@ namespace rocwmma using TestBlockSizes32x32LargeBlockK = std::tuple, I<32>, I<8>>, std::tuple, I<32>, I<16>>, std::tuple, I<32>, I<32>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<32>, I<64>> #endif // ROCWMMA_EXTENDED_TESTS @@ -246,7 +246,7 @@ namespace rocwmma using TestBlockSizes32x32HugeBlockK = std::tuple, I<32>, I<8>>, std::tuple, I<32>, I<16>>, std::tuple, I<32>, I<32>> -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<32>, I<64>>, std::tuple, I<32>, I<128>> @@ -298,10 +298,10 @@ namespace rocwmma { // clang-format off // Don't benchmark wg less than 4 waves by default -#if defined(ROCWMMA_VALIDATION_TESTS) || defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_VALIDATION_TESTS || ROCWMMA_EXTENDED_TESTS {warpSize, 1}, // 1 wave {warpSize, 2}, {warpSize * 2, 1}, // 2 wave -#endif // ROCWMMA_VALIDATION_TESTS +#endif // ROCWMMA_VALIDATION_TESTS || ROCWMMA_EXTENDED_TESTS {warpSize, 4}, {warpSize * 2, 2}, // 4 wave {warpSize * 4, 1} // 4 wave // clang-format on @@ -310,7 +310,6 @@ namespace rocwmma static inline std::vector problemSizes() { - return { // clang-format off @@ -323,7 +322,7 @@ namespace rocwmma {512, 512, 512}, // Skip validation on larger sizes // due to very slow. -#if !defined(ROCWMMA_VALIDATION_TESTS) +#if !ROCWMMA_VALIDATION_TESTS {1024, 1024, 1024}, {2048, 2048, 2048}, {2560, 2560, 2560}, @@ -331,7 +330,7 @@ namespace rocwmma {3584, 3584, 3584}, {4096, 4096, 4096}, {5120, 5120, 5120}, -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS {6144, 6144, 6144}, {7168, 7168, 7168}, {8192, 8192, 8192}, diff --git a/test/gemm/gemm_kernel_base.cpp b/test/gemm/gemm_kernel_base.cpp index 1292e3bb..744aad24 100644 --- a/test/gemm/gemm_kernel_base.cpp +++ b/test/gemm/gemm_kernel_base.cpp @@ -845,7 +845,7 @@ namespace rocwmma ROCWMMA_INSTANTIATE_GEMM_KERNEL_BASE(xfloat32_t, float32_t, float32_t); ROCWMMA_INSTANTIATE_GEMM_KERNEL_BASE(float64_t, float64_t, float64_t); -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS ROCWMMA_INSTANTIATE_GEMM_KERNEL_BASE(int8_t, int8_t, int32_t); ROCWMMA_INSTANTIATE_GEMM_KERNEL_BASE(bfloat16_t, bfloat16_t, bfloat16_t); ROCWMMA_INSTANTIATE_GEMM_KERNEL_BASE(bfloat16_t, bfloat16_t, float32_t); diff --git a/test/gemm/gemm_kernel_base_impl.hpp b/test/gemm/gemm_kernel_base_impl.hpp index 937838a7..44849096 100644 --- a/test/gemm/gemm_kernel_base_impl.hpp +++ b/test/gemm/gemm_kernel_base_impl.hpp @@ -49,7 +49,7 @@ #if ROCWMMA_ROCBLAS_INTEGRATION #include "rocblas_reference.hpp" // rocBLAS GPU kernel -#endif // ROCWMMA_VALIDATE_WITH_ROCBLAS || ROCWMMA_BENCHMARK_WITH_ROCBLAS +#endif // ROCWMMA_ROCBLAS_INTEGRATION namespace rocwmma { @@ -496,7 +496,7 @@ namespace rocwmma if(!mRunFlag) { - stream << "n/a, " + stream << "n/a" << ", " << "n/a" << ", " diff --git a/test/gemm/gemm_resource.cpp b/test/gemm/gemm_resource.cpp index 42e4d04c..714d82ea 100644 --- a/test/gemm/gemm_resource.cpp +++ b/test/gemm/gemm_resource.cpp @@ -40,7 +40,7 @@ namespace rocwmma template struct GemmResource; template struct GemmResource; -#if defined(ROCWMMA_EXTENDED_TESTS) +#if ROCWMMA_EXTENDED_TESTS template struct GemmResource; template struct GemmResource; template struct GemmResource; diff --git a/test/hip_device.cpp b/test/hip_device.cpp index 420f2c7f..5ce1364b 100644 --- a/test/hip_device.cpp +++ b/test/hip_device.cpp @@ -91,7 +91,7 @@ namespace rocwmma mMaxFreqMhz = static_cast(static_cast(mProps.clockRate) / 1000.0); mCurFreqMhz = mMaxFreqMhz; -#ifdef ROCWMMA_BENCHMARK_TESTS +#if ROCWMMA_BENCHMARK_TESTS bool smiErrorFlag = false; CHECK_RSMI_ERROR(rsmi_init(0), smiErrorFlag); if(!smiErrorFlag) @@ -134,7 +134,7 @@ namespace rocwmma } } } -#endif +#endif // ROCWMMA_BENCHMARK_TESTS } hipDevice_t HipDevice::getDeviceHandle() const @@ -184,10 +184,10 @@ namespace rocwmma HipDevice::~HipDevice() { -#ifdef ROCWMMA_BENCHMARK_TESTS +#if ROCWMMA_BENCHMARK_TESTS bool smiErrorFlag = false; CHECK_RSMI_ERROR(rsmi_shut_down(), smiErrorFlag); -#endif +#endif // ROCWMMA_BENCHMARK_TESTS } // Need to check the host device target support statically before hip modules attempt diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_128.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_128.cpp index c427b4fe..20462fde 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_128.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_128.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_16.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_16.cpp index a7ae76dc..3a62d5aa 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_16.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_16.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_256.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_256.cpp index 5d37441f..586bdb0a 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_256.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_256.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_32.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_32.cpp index 540e56e3..89cb62d1 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_32.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_32.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_64.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_64.cpp index e955dbe1..b68fc818 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_64.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_a_64.cpp @@ -67,17 +67,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_128.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_128.cpp index 0542829e..d7957ca7 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_128.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_128.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_16.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_16.cpp index d256fadc..733c4557 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_16.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_16.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_256.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_256.cpp index 34389874..17f2a155 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_256.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_256.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_32.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_32.cpp index fe18c36c..78ab518e 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_32.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_32.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_64.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_64.cpp index 143da6ee..529bc84f 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_64.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_acc_64.cpp @@ -67,17 +67,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_128.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_128.cpp index 3adda405..58edf755 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_128.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_128.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_16.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_16.cpp index 71a7b4f0..61e87b51 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_16.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_16.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_256.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_256.cpp index 272943a0..4537cf73 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_256.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_256.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_32.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_32.cpp index 46cf446f..90410bb4 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_32.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_32.cpp @@ -66,17 +66,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_64.cpp b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_64.cpp index 6d73dafe..15beeaaa 100644 --- a/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_64.cpp +++ b/test/unit/load_store_matrix_coop_sync_test/test/load_store_matrix_coop_sync_b_64.cpp @@ -67,17 +67,14 @@ namespace rocwmma static inline std::vector param2s() { - return { - 0.0, - 1.0, - 2.0, - 3.0 // 1 - 4 waves -#ifdef ROCWMMA_EXTENDED_TESTS - , - 4.0, - 5.0, - 6.0, - 7.0 // 8 waves + return + { + 0.0, 1.0, 2.0, + 3.0 // 1 - 4 waves +#if ROCWMMA_EXTENDED_TESTS + , + 4.0, 5.0, 6.0, + 7.0 // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; } diff --git a/test/unit/unit_test_params.hpp b/test/unit/unit_test_params.hpp index 25a7d735..5b508763 100644 --- a/test/unit/unit_test_params.hpp +++ b/test/unit/unit_test_params.hpp @@ -58,7 +58,7 @@ namespace rocwmma xfloat32_t, int8_t, int32_t -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS , uint8_t, uint32_t @@ -81,7 +81,7 @@ namespace rocwmma using TestBlockSizes16 = std::tuple, I<16>>, std::tuple, I<32>>, std::tuple, I<64>> -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<128>>, std::tuple, I<256>> @@ -92,7 +92,7 @@ namespace rocwmma std::tuple, I<16>>, std::tuple, I<32>>, std::tuple, I<64>> -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<128>>, std::tuple, I<256>> @@ -103,7 +103,7 @@ namespace rocwmma std::tuple, I<16>>, std::tuple, I<32>>, std::tuple, I<64>> -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<128>>, std::tuple, I<256>> @@ -114,7 +114,7 @@ namespace rocwmma std::tuple, I<16>>, std::tuple, I<32>>, std::tuple, I<64>> -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<128>>, std::tuple, I<256>> @@ -125,7 +125,7 @@ namespace rocwmma std::tuple, I<16>>, std::tuple, I<32>>, std::tuple, I<64>> -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS , std::tuple, I<128>>, std::tuple, I<256>> @@ -156,7 +156,7 @@ namespace rocwmma return { {warpSize, 1}, // 1 Wave {warpSize, 2}, {warpSize * 2, 1}, // 2 Waves {warpSize, 4}, {warpSize * 2, 2}, {warpSize * 4, 1}, // 4 Waves -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS {warpSize, 8}, {warpSize * 2, 4}, {warpSize * 4, 2}, {warpSize * 8, 1} // 8 waves #endif // ROCWMMA_EXTENDED_TESTS }; @@ -175,7 +175,7 @@ namespace rocwmma {512, 8}, {512, 16}, {512, 32}, {512, 64}, {512, 128}, {512, 256}, {512, 512}, {1024, 8}, {1024, 16}, {1024, 32}, {1024, 64}, {1024, 128}, {1024, 256}, {1024, 512}, -#ifdef ROCWMMA_EXTENDED_TESTS +#if ROCWMMA_EXTENDED_TESTS {1024, 1024}, {2048, 2048}, {2560, 2560},