diff --git a/Examples/AtfCCSD/AtfCCSD.cpp b/Examples/AtfCCSD/AtfCCSD.cpp index c2388210..a4b1e4a8 100644 --- a/Examples/AtfCCSD/AtfCCSD.cpp +++ b/Examples/AtfCCSD/AtfCCSD.cpp @@ -5,9 +5,9 @@ using namespace std; class AtfCCSD : public ExampleBase { protected: - AtfCCSD(shared_ptr config, int defaultProblemSize, + AtfCCSD(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleBase(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) + ExampleBase(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) { // Keep OpenCL sizes as specified m_inputSize1 = 24; @@ -18,7 +18,7 @@ class AtfCCSD : public ExampleBase { m_inputSize6 = 16; m_inputSize7 = 24; - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); + m_tuner->SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); const std::string atfSamplesPath1 = GetKernelFilePath(exampleFolderPath, "TcAbcdefGebcDfga1"); const std::string atfSamplesPath2 = GetKernelFilePath(exampleFolderPath, "TcAbcdefGebcDfga2"); @@ -105,19 +105,19 @@ class AtfCCSD : public ExampleBase { void InitKernel() override { - m_aId = m_tuner.AddArgumentVector(m_a, ktt::ArgumentAccessType::ReadOnly); - m_bId = m_tuner.AddArgumentVector(m_b, ktt::ArgumentAccessType::ReadOnly); - m_cId = m_tuner.AddArgumentVector(m_c, ktt::ArgumentAccessType::ReadWrite); - m_intResId = m_tuner.AddArgumentVector(m_intRes, ktt::ArgumentAccessType::ReadWrite); - m_resId = m_tuner.AddArgumentVector(m_res, ktt::ArgumentAccessType::ReadWrite); + m_aId = m_tuner->AddArgumentVector(m_a, ktt::ArgumentAccessType::ReadOnly); + m_bId = m_tuner->AddArgumentVector(m_b, ktt::ArgumentAccessType::ReadOnly); + m_cId = m_tuner->AddArgumentVector(m_c, ktt::ArgumentAccessType::ReadWrite); + m_intResId = m_tuner->AddArgumentVector(m_intRes, ktt::ArgumentAccessType::ReadWrite); + m_resId = m_tuner->AddArgumentVector(m_res, ktt::ArgumentAccessType::ReadWrite); - m_definition = m_tuner.AddKernelDefinitionFromFile("tc_1", m_kernelPath1, ktt::DimensionVector(), ktt::DimensionVector()); - m_definition2 = m_tuner.AddKernelDefinitionFromFile("tc_2", m_kernelPath2, ktt::DimensionVector(), ktt::DimensionVector()); + m_definition = m_tuner->AddKernelDefinitionFromFile("tc_1", m_kernelPath1, ktt::DimensionVector(), ktt::DimensionVector()); + m_definition2 = m_tuner->AddKernelDefinitionFromFile("tc_2", m_kernelPath2, ktt::DimensionVector(), ktt::DimensionVector()); - m_tuner.SetArguments(m_definition, {m_aId, m_bId, m_resId, m_intResId}); - m_tuner.SetArguments(m_definition2, {m_intResId, m_resId, m_cId}); + m_tuner->SetArguments(m_definition, {m_aId, m_bId, m_resId, m_intResId}); + m_tuner->SetArguments(m_definition2, {m_intResId, m_resId, m_cId}); - m_kernel = m_tuner.CreateCompositeKernel("CCSD", {m_definition, m_definition2}, [this](ktt::ComputeInterface& interface) + m_kernel = m_tuner->CreateCompositeKernel("CCSD", {m_definition, m_definition2}, [this](ktt::ComputeInterface& interface) { const auto& pairs = interface.GetCurrentConfiguration().GetPairs(); size_t newResSize = m_resSize; @@ -193,135 +193,135 @@ class AtfCCSD : public ExampleBase { auto NoPostInSecondKernelConstraint = [](const vector& v) { return v[0] == 1 || (v[0] % v[1] == 0); }; // Add parameters - m_tuner.AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); - m_tuner.AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - m_tuner.AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1, 2, 3, 4, 5, 6}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_2", vector{0, 1, 2, 3, 4, 5, 6}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_3", vector{0, 1, 2, 3, 4, 5, 6}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_4", vector{0, 1, 2, 3, 4, 5, 6}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_5", vector{0, 1, 2, 3, 4, 5, 6}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_6", vector{0, 1, 2, 3, 4, 5, 6}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_R_1", vector{0, 1, 2, 3, 4, 5, 6}); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_2", vector{m_inputSize2}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_2", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_2", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_2", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_2", ParameterRange(m_inputSize2)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_3", vector{m_inputSize3}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_3", ParameterRange(m_inputSize3)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_3", ParameterRange(m_inputSize3)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_3", ParameterRange(m_inputSize3)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_3", ParameterRange(m_inputSize3)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_4", vector{m_inputSize4}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_4", ParameterRange(m_inputSize4)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_4", ParameterRange(m_inputSize4)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_4", ParameterRange(m_inputSize4)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_4", ParameterRange(m_inputSize4)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_5", vector{m_inputSize5}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_5", ParameterRange(m_inputSize5)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_5", ParameterRange(m_inputSize5)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_5", ParameterRange(m_inputSize5)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_5", ParameterRange(m_inputSize5)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_6", vector{m_inputSize6}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_6", ParameterRange(m_inputSize6)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_6", ParameterRange(m_inputSize6)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_6", ParameterRange(m_inputSize6)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_6", ParameterRange(m_inputSize6)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_R_1", vector{m_inputSize7}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_R_1", ParameterRange(m_inputSize7)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_R_1", ParameterRange(m_inputSize7)); - m_tuner.AddParameter(m_kernel, "NUM_WG_R_1", ParameterRange(m_inputSize7)); - m_tuner.AddParameter(m_kernel, "NUM_WI_R_1", ParameterRange(m_inputSize7)); - - m_tuner.AddParameter(m_kernel, "L_REDUCTION", vector{1}); - m_tuner.AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); - m_tuner.AddParameter(m_kernel, "L_WRITE_BACK", vector{6}); + m_tuner->AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); + m_tuner->AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + m_tuner->AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1, 2, 3, 4, 5, 6}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_2", vector{0, 1, 2, 3, 4, 5, 6}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_3", vector{0, 1, 2, 3, 4, 5, 6}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_4", vector{0, 1, 2, 3, 4, 5, 6}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_5", vector{0, 1, 2, 3, 4, 5, 6}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_6", vector{0, 1, 2, 3, 4, 5, 6}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_R_1", vector{0, 1, 2, 3, 4, 5, 6}); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_2", vector{m_inputSize2}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_2", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_2", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_2", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_2", ParameterRange(m_inputSize2)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_3", vector{m_inputSize3}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_3", ParameterRange(m_inputSize3)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_3", ParameterRange(m_inputSize3)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_3", ParameterRange(m_inputSize3)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_3", ParameterRange(m_inputSize3)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_4", vector{m_inputSize4}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_4", ParameterRange(m_inputSize4)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_4", ParameterRange(m_inputSize4)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_4", ParameterRange(m_inputSize4)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_4", ParameterRange(m_inputSize4)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_5", vector{m_inputSize5}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_5", ParameterRange(m_inputSize5)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_5", ParameterRange(m_inputSize5)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_5", ParameterRange(m_inputSize5)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_5", ParameterRange(m_inputSize5)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_6", vector{m_inputSize6}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_6", ParameterRange(m_inputSize6)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_6", ParameterRange(m_inputSize6)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_6", ParameterRange(m_inputSize6)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_6", ParameterRange(m_inputSize6)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_R_1", vector{m_inputSize7}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_R_1", ParameterRange(m_inputSize7)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_R_1", ParameterRange(m_inputSize7)); + m_tuner->AddParameter(m_kernel, "NUM_WG_R_1", ParameterRange(m_inputSize7)); + m_tuner->AddParameter(m_kernel, "NUM_WI_R_1", ParameterRange(m_inputSize7)); + + m_tuner->AddParameter(m_kernel, "L_REDUCTION", vector{1}); + m_tuner->AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); + m_tuner->AddParameter(m_kernel, "L_WRITE_BACK", vector{6}); // Add constraints - m_tuner.AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_2"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_3"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_4"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_5"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_6"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_R_1"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_3"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_4"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_5"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_6"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_R_1"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_L_4"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_L_5"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_L_6"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_R_1"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_4", "OCL_DIM_L_5"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_4", "OCL_DIM_L_6"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_4", "OCL_DIM_R_1"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_5", "OCL_DIM_L_6"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_5", "OCL_DIM_R_1"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_6", "OCL_DIM_R_1"}, UnequalConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_2", "INPUT_SIZE_L_2"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_2", "INPUT_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_2", "L_CB_SIZE_L_2", "P_CB_SIZE_L_2"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_2", "INPUT_SIZE_L_2", "NUM_WG_L_2"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_3", "INPUT_SIZE_L_3"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_3", "L_CB_SIZE_L_3"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_3", "INPUT_SIZE_L_3", "L_CB_SIZE_L_3"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_3", "L_CB_SIZE_L_3", "P_CB_SIZE_L_3"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_3", "INPUT_SIZE_L_3", "NUM_WG_L_3"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_4", "INPUT_SIZE_L_4"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_4", "L_CB_SIZE_L_4"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_4", "INPUT_SIZE_L_4", "L_CB_SIZE_L_4"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_4", "L_CB_SIZE_L_4", "P_CB_SIZE_L_4"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_4", "INPUT_SIZE_L_4", "NUM_WG_L_4"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_5", "INPUT_SIZE_L_5"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_5", "L_CB_SIZE_L_5"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_5", "INPUT_SIZE_L_5", "L_CB_SIZE_L_5"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_5", "L_CB_SIZE_L_5", "P_CB_SIZE_L_5"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_5", "INPUT_SIZE_L_5", "NUM_WG_L_5"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_6", "INPUT_SIZE_L_6"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_6", "L_CB_SIZE_L_6"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_6", "INPUT_SIZE_L_6", "L_CB_SIZE_L_6"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_6", "L_CB_SIZE_L_6", "P_CB_SIZE_L_6"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_6", "INPUT_SIZE_L_6", "NUM_WG_L_6"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_R_1", "INPUT_SIZE_R_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_R_1", "INPUT_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_R_1", "L_CB_SIZE_R_1", "P_CB_SIZE_R_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_R_1", "INPUT_SIZE_R_1", "NUM_WG_R_1"}, LessThanOrEqualCeilDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_R_1", "L_CB_SIZE_R_1"}, NoPostInSecondKernelConstraint); + m_tuner->AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_2"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_3"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_4"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_5"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_6"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_R_1"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_3"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_4"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_5"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_L_6"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_2", "OCL_DIM_R_1"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_L_4"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_L_5"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_L_6"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_3", "OCL_DIM_R_1"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_4", "OCL_DIM_L_5"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_4", "OCL_DIM_L_6"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_4", "OCL_DIM_R_1"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_5", "OCL_DIM_L_6"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_5", "OCL_DIM_R_1"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_6", "OCL_DIM_R_1"}, UnequalConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_2", "INPUT_SIZE_L_2"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_2", "INPUT_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_2", "L_CB_SIZE_L_2", "P_CB_SIZE_L_2"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_2", "INPUT_SIZE_L_2", "NUM_WG_L_2"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_3", "INPUT_SIZE_L_3"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_3", "L_CB_SIZE_L_3"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_3", "INPUT_SIZE_L_3", "L_CB_SIZE_L_3"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_3", "L_CB_SIZE_L_3", "P_CB_SIZE_L_3"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_3", "INPUT_SIZE_L_3", "NUM_WG_L_3"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_4", "INPUT_SIZE_L_4"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_4", "L_CB_SIZE_L_4"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_4", "INPUT_SIZE_L_4", "L_CB_SIZE_L_4"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_4", "L_CB_SIZE_L_4", "P_CB_SIZE_L_4"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_4", "INPUT_SIZE_L_4", "NUM_WG_L_4"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_5", "INPUT_SIZE_L_5"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_5", "L_CB_SIZE_L_5"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_5", "INPUT_SIZE_L_5", "L_CB_SIZE_L_5"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_5", "L_CB_SIZE_L_5", "P_CB_SIZE_L_5"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_5", "INPUT_SIZE_L_5", "NUM_WG_L_5"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_6", "INPUT_SIZE_L_6"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_6", "L_CB_SIZE_L_6"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_6", "INPUT_SIZE_L_6", "L_CB_SIZE_L_6"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_6", "L_CB_SIZE_L_6", "P_CB_SIZE_L_6"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_6", "INPUT_SIZE_L_6", "NUM_WG_L_6"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_R_1", "INPUT_SIZE_R_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_R_1", "INPUT_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_R_1", "L_CB_SIZE_R_1", "P_CB_SIZE_R_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_R_1", "INPUT_SIZE_R_1", "NUM_WG_R_1"}, LessThanOrEqualCeilDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_R_1", "L_CB_SIZE_R_1"}, NoPostInSecondKernelConstraint); // Thread modifiers for first kernel (definition) - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WG_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WG_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WG_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WG_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WG_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -335,7 +335,7 @@ class AtfCCSD : public ExampleBase { + static_cast(values[18] == 0) * values[19] * values[20]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WG_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WG_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WG_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WG_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WG_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -349,7 +349,7 @@ class AtfCCSD : public ExampleBase { + static_cast(values[18] == 1) * values[19] * values[20]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WG_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WG_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WG_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WG_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WG_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -364,7 +364,7 @@ class AtfCCSD : public ExampleBase { }); // Thread modifiers for second kernel (definition2) - m_tuner.AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WG_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WG_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WG_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WG_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -378,7 +378,7 @@ class AtfCCSD : public ExampleBase { + static_cast(values[18] == 0) * values[19]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WG_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WG_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WG_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WG_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -392,7 +392,7 @@ class AtfCCSD : public ExampleBase { + static_cast(values[18] == 1) * values[19]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WG_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WG_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WG_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WG_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -407,7 +407,7 @@ class AtfCCSD : public ExampleBase { }); // Local thread modifiers - m_tuner.AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -421,7 +421,7 @@ class AtfCCSD : public ExampleBase { + static_cast(values[12] == 0) * values[13]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) @@ -435,7 +435,7 @@ class AtfCCSD : public ExampleBase { + static_cast(values[12] == 1) * values[13]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Z, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2", "OCL_DIM_L_3", "NUM_WI_L_3", "OCL_DIM_L_4", "NUM_WI_L_4", "OCL_DIM_L_5", "NUM_WI_L_5", "OCL_DIM_L_6", "NUM_WI_L_6", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) diff --git a/Examples/AtfConvolution/AtfConvolution.cpp b/Examples/AtfConvolution/AtfConvolution.cpp index 493f4206..39358346 100644 --- a/Examples/AtfConvolution/AtfConvolution.cpp +++ b/Examples/AtfConvolution/AtfConvolution.cpp @@ -5,15 +5,15 @@ using namespace std; class AtfConvolution : public ExampleBase { protected: - AtfConvolution(shared_ptr config, int defaultProblemSize, + AtfConvolution(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleBase(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) + ExampleBase(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) { // Keep OpenCL sizes as specified m_inputSize1 = static_cast(sqrt(m_problemSize)) * 1024; m_inputSize2 = m_inputSize1; - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); + m_tuner->SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); } friend ExampleBase; @@ -70,9 +70,9 @@ class AtfConvolution : public ExampleBase { void InitKernel() override { - m_inId = m_tuner.AddArgumentVector(m_in, ktt::ArgumentAccessType::ReadOnly); - m_outId = m_tuner.AddArgumentVector(m_out, ktt::ArgumentAccessType::ReadWrite); - m_intResId = m_tuner.AddArgumentVector(m_intRes, ktt::ArgumentAccessType::ReadWrite); + m_inId = m_tuner->AddArgumentVector(m_in, ktt::ArgumentAccessType::ReadOnly); + m_outId = m_tuner->AddArgumentVector(m_out, ktt::ArgumentAccessType::ReadWrite); + m_intResId = m_tuner->AddArgumentVector(m_intRes, ktt::ArgumentAccessType::ReadWrite); InitKernelDefault("gaussian_1", "Convolution", ktt::DimensionVector(), {m_inId, m_outId, m_intResId}); } @@ -115,49 +115,49 @@ class AtfConvolution : public ExampleBase { auto DividesDivConstraint = [](const vector& v) { return (v[1] / v[2]) % v[0] == 0; }; // Add parameters - m_tuner.AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); - m_tuner.AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - m_tuner.AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_2", vector{0, 1}); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1 - 4}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1 - 4)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1 - 4)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1 - 4)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1 - 4)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_2", vector{m_inputSize2 - 4}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_2", ParameterRange(m_inputSize2 - 4)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_2", ParameterRange(m_inputSize2 - 4)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_2", ParameterRange(m_inputSize2 - 4)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_2", ParameterRange(m_inputSize2 - 4)); - - m_tuner.AddParameter(m_kernel, "L_REDUCTION", vector{1}); - m_tuner.AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); - m_tuner.AddParameter(m_kernel, "L_WRITE_BACK", vector{2}); + m_tuner->AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); + m_tuner->AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + m_tuner->AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_2", vector{0, 1}); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1 - 4}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1 - 4)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1 - 4)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1 - 4)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1 - 4)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_2", vector{m_inputSize2 - 4}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_2", ParameterRange(m_inputSize2 - 4)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_2", ParameterRange(m_inputSize2 - 4)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_2", ParameterRange(m_inputSize2 - 4)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_2", ParameterRange(m_inputSize2 - 4)); + + m_tuner->AddParameter(m_kernel, "L_REDUCTION", vector{1}); + m_tuner->AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); + m_tuner->AddParameter(m_kernel, "L_WRITE_BACK", vector{2}); // Add constraints - m_tuner.AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_2"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_2"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_2", "INPUT_SIZE_L_2"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_2", "INPUT_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_2", "L_CB_SIZE_L_2", "P_CB_SIZE_L_2"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_2", "INPUT_SIZE_L_2", "NUM_WG_L_2"}, LessThanOrEqualCeilDivConstraint); + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_2", "INPUT_SIZE_L_2"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_2", "INPUT_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_2", "L_CB_SIZE_L_2", "P_CB_SIZE_L_2"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_2", "INPUT_SIZE_L_2", "NUM_WG_L_2"}, LessThanOrEqualCeilDivConstraint); // Thread modifiers for global X dimension - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2"}, [](const uint64_t, const vector& values) { @@ -166,7 +166,7 @@ class AtfConvolution : public ExampleBase { }); // Thread modifiers for global Y dimension - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2"}, [](const uint64_t, const vector& values) { @@ -175,7 +175,7 @@ class AtfConvolution : public ExampleBase { }); // Thread modifiers for local X dimension - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2"}, [](const uint64_t, const vector& values) { @@ -184,7 +184,7 @@ class AtfConvolution : public ExampleBase { }); // Thread modifiers for local Y dimension - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2"}, [](const uint64_t, const vector& values) { diff --git a/Examples/AtfGEMM/AtfGEMM.cpp b/Examples/AtfGEMM/AtfGEMM.cpp index 639ef0be..2ce688a6 100644 --- a/Examples/AtfGEMM/AtfGEMM.cpp +++ b/Examples/AtfGEMM/AtfGEMM.cpp @@ -5,16 +5,16 @@ using namespace std; class AtfGEMM : public ExampleBase { protected: - AtfGEMM(shared_ptr config, int defaultProblemSize, + AtfGEMM(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleBase(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) + ExampleBase(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) { // Keep OpenCL sizes as specified m_inputSize1 = static_cast(sqrt(m_problemSize)) * 1024; m_inputSize2 = m_inputSize1; m_inputSize3 = m_inputSize1; - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); + m_tuner->SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); m_kernelPath1 = GetKernelFilePath(exampleFolderPath, "Gemm1"); m_kernelPath2 = GetKernelFilePath(exampleFolderPath, "Gemm2"); @@ -95,19 +95,19 @@ class AtfGEMM : public ExampleBase { void InitKernel() override { - m_aId = m_tuner.AddArgumentVector(m_a, ktt::ArgumentAccessType::ReadOnly); - m_bId = m_tuner.AddArgumentVector(m_b, ktt::ArgumentAccessType::ReadOnly); - m_cId = m_tuner.AddArgumentVector(m_c, ktt::ArgumentAccessType::ReadWrite); - m_intResId = m_tuner.AddArgumentVector(m_intRes, ktt::ArgumentAccessType::ReadWrite); - m_resId = m_tuner.AddArgumentVector(m_res, ktt::ArgumentAccessType::ReadWrite); + m_aId = m_tuner->AddArgumentVector(m_a, ktt::ArgumentAccessType::ReadOnly); + m_bId = m_tuner->AddArgumentVector(m_b, ktt::ArgumentAccessType::ReadOnly); + m_cId = m_tuner->AddArgumentVector(m_c, ktt::ArgumentAccessType::ReadWrite); + m_intResId = m_tuner->AddArgumentVector(m_intRes, ktt::ArgumentAccessType::ReadWrite); + m_resId = m_tuner->AddArgumentVector(m_res, ktt::ArgumentAccessType::ReadWrite); - m_definition = m_tuner.AddKernelDefinitionFromFile("gemm_1", m_kernelPath1, ktt::DimensionVector(), ktt::DimensionVector()); - m_definition2 = m_tuner.AddKernelDefinitionFromFile("gemm_2", m_kernelPath2, ktt::DimensionVector(), ktt::DimensionVector()); + m_definition = m_tuner->AddKernelDefinitionFromFile("gemm_1", m_kernelPath1, ktt::DimensionVector(), ktt::DimensionVector()); + m_definition2 = m_tuner->AddKernelDefinitionFromFile("gemm_2", m_kernelPath2, ktt::DimensionVector(), ktt::DimensionVector()); - m_tuner.SetArguments(m_definition, {m_aId, m_bId, m_resId, m_intResId}); - m_tuner.SetArguments(m_definition2, {m_intResId, m_resId, m_cId}); + m_tuner->SetArguments(m_definition, {m_aId, m_bId, m_resId, m_intResId}); + m_tuner->SetArguments(m_definition2, {m_intResId, m_resId, m_cId}); - m_kernel = m_tuner.CreateCompositeKernel("GEMM", {m_definition, m_definition2}, [this](ktt::ComputeInterface& interface) + m_kernel = m_tuner->CreateCompositeKernel("GEMM", {m_definition, m_definition2}, [this](ktt::ComputeInterface& interface) { const auto& pairs = interface.GetCurrentConfiguration().GetPairs(); size_t newResSize = m_resSize; @@ -183,65 +183,65 @@ class AtfGEMM : public ExampleBase { auto NoPostInSecondKernelConstraint = [](const vector& v) { return v[0] == 1 || (v[0] % v[1] == 0); }; // Add parameters - m_tuner.AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); - m_tuner.AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - m_tuner.AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1, 2}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_2", vector{0, 1, 2}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_R_1", vector{0, 1, 2}); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_2", vector{m_inputSize2}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_2", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_2", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_2", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_2", ParameterRange(m_inputSize2)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_R_1", vector{m_inputSize3}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_R_1", ParameterRange(m_inputSize3)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_R_1", ParameterRange(m_inputSize3)); - m_tuner.AddParameter(m_kernel, "NUM_WG_R_1", ParameterRange(m_inputSize3)); - m_tuner.AddParameter(m_kernel, "NUM_WI_R_1", ParameterRange(m_inputSize3)); - - m_tuner.AddParameter(m_kernel, "L_REDUCTION", vector{1}); - m_tuner.AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); - m_tuner.AddParameter(m_kernel, "L_WRITE_BACK", vector{2}); + m_tuner->AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); + m_tuner->AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + m_tuner->AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1, 2}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_2", vector{0, 1, 2}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_R_1", vector{0, 1, 2}); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_2", vector{m_inputSize2}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_2", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_2", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_2", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_2", ParameterRange(m_inputSize2)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_R_1", vector{m_inputSize3}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_R_1", ParameterRange(m_inputSize3)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_R_1", ParameterRange(m_inputSize3)); + m_tuner->AddParameter(m_kernel, "NUM_WG_R_1", ParameterRange(m_inputSize3)); + m_tuner->AddParameter(m_kernel, "NUM_WI_R_1", ParameterRange(m_inputSize3)); + + m_tuner->AddParameter(m_kernel, "L_REDUCTION", vector{1}); + m_tuner->AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); + m_tuner->AddParameter(m_kernel, "L_WRITE_BACK", vector{2}); // Add constraints - m_tuner.AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_2"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_R_1", "OCL_DIM_L_2"}, UnequalConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_R_1"}, UnequalConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_2", "INPUT_SIZE_L_2"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_2", "INPUT_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_2", "L_CB_SIZE_L_2", "P_CB_SIZE_L_2"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_2", "INPUT_SIZE_L_2", "NUM_WG_L_2"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_R_1", "INPUT_SIZE_R_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_R_1", "INPUT_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_R_1", "L_CB_SIZE_R_1", "P_CB_SIZE_R_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_R_1", "INPUT_SIZE_R_1", "NUM_WG_R_1"}, LessThanOrEqualCeilDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_R_1", "L_CB_SIZE_R_1"}, NoPostInSecondKernelConstraint); + m_tuner->AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_L_2"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_R_1", "OCL_DIM_L_2"}, UnequalConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_R_1"}, UnequalConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_2", "INPUT_SIZE_L_2"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_2", "INPUT_SIZE_L_2", "L_CB_SIZE_L_2"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_2", "L_CB_SIZE_L_2", "P_CB_SIZE_L_2"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_2", "INPUT_SIZE_L_2", "NUM_WG_L_2"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_R_1", "INPUT_SIZE_R_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_R_1", "INPUT_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_R_1", "L_CB_SIZE_R_1", "P_CB_SIZE_R_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_R_1", "INPUT_SIZE_R_1", "NUM_WG_R_1"}, LessThanOrEqualCeilDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_R_1", "L_CB_SIZE_R_1"}, NoPostInSecondKernelConstraint); // Thread modifiers for first kernel (definition) - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WG_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -250,7 +250,7 @@ class AtfGEMM : public ExampleBase { + static_cast(values[6] == 0) * values[7] * values[8]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WG_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -259,7 +259,7 @@ class AtfGEMM : public ExampleBase { + static_cast(values[6] == 1) * values[7] * values[8]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WG_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -269,7 +269,7 @@ class AtfGEMM : public ExampleBase { }); // Thread modifiers for second kernel (definition2) - m_tuner.AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -278,7 +278,7 @@ class AtfGEMM : public ExampleBase { + static_cast(values[6] == 0) * values[7]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -287,7 +287,7 @@ class AtfGEMM : public ExampleBase { + static_cast(values[6] == 1) * values[7]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_definition2}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, {"OCL_DIM_L_1", "NUM_WG_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WG_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -297,7 +297,7 @@ class AtfGEMM : public ExampleBase { }); // Local thread modifiers - m_tuner.AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::X, + m_tuner->AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -306,7 +306,7 @@ class AtfGEMM : public ExampleBase { + static_cast(values[4] == 0) * values[5]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, + m_tuner->AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { @@ -315,7 +315,7 @@ class AtfGEMM : public ExampleBase { + static_cast(values[4] == 1) * values[5]; }); - m_tuner.AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_definition, m_definition2}, ktt::ModifierType::Local, ktt::ModifierDimension::Z, {"OCL_DIM_L_1", "NUM_WI_L_1", "OCL_DIM_L_2", "NUM_WI_L_2", "OCL_DIM_R_1", "NUM_WI_R_1"}, [](const uint64_t, const vector& values) { diff --git a/Examples/AtfPRL/AtfPRL.cpp b/Examples/AtfPRL/AtfPRL.cpp index a0923163..3ef8e5a9 100644 --- a/Examples/AtfPRL/AtfPRL.cpp +++ b/Examples/AtfPRL/AtfPRL.cpp @@ -6,15 +6,15 @@ using namespace std; class AtfPRL : public ExampleBase { protected: - AtfPRL(shared_ptr config, int defaultProblemSize, + AtfPRL(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleBase(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) + ExampleBase(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) { // Keep OpenCL sizes as specified m_inputSize1 = static_cast(sqrt(m_problemSize)) * 1024; m_inputSize2 = m_inputSize1; - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); + m_tuner->SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); } friend ExampleBase; @@ -63,8 +63,8 @@ class AtfPRL : public ExampleBase { void InitKernel() override { - m_aId = m_tuner.AddArgumentVector(m_a, ktt::ArgumentAccessType::ReadOnly); - m_bId = m_tuner.AddArgumentVector(m_b, ktt::ArgumentAccessType::ReadWrite); + m_aId = m_tuner->AddArgumentVector(m_a, ktt::ArgumentAccessType::ReadOnly); + m_bId = m_tuner->AddArgumentVector(m_b, ktt::ArgumentAccessType::ReadWrite); InitKernelDefault("rl_1", "PRL", ktt::DimensionVector(), {m_aId, m_bId}); } @@ -108,47 +108,47 @@ class AtfPRL : public ExampleBase { auto NoPostInSecondKernelConstraint = [](const vector& v) { return v[0] == 1 || (v[0] % v[1] == 0); }; // Add parameters - m_tuner.AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); - m_tuner.AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - m_tuner.AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); - - m_tuner.AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "OCL_DIM_R_1", vector{0, 1}); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1)); - m_tuner.AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1)); - - m_tuner.AddParameter(m_kernel, "INPUT_SIZE_R_1", vector{m_inputSize2}); - m_tuner.AddParameter(m_kernel, "L_CB_SIZE_R_1", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "P_CB_SIZE_R_1", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "NUM_WG_R_1", ParameterRange(m_inputSize2)); - m_tuner.AddParameter(m_kernel, "NUM_WI_R_1", ParameterRange(m_inputSize2)); - - m_tuner.AddParameter(m_kernel, "L_REDUCTION", vector{1}); - m_tuner.AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); - m_tuner.AddParameter(m_kernel, "L_WRITE_BACK", vector{1}); + m_tuner->AddParameter(m_kernel, "CACHE_L_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "CACHE_P_CB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "G_CB_RES_DEST_LEVEL", vector{2}); + m_tuner->AddParameter(m_kernel, "L_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + m_tuner->AddParameter(m_kernel, "P_CB_RES_DEST_LEVEL", vector{2, 1, 0}); + + m_tuner->AddParameter(m_kernel, "OCL_DIM_L_1", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "OCL_DIM_R_1", vector{0, 1}); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_L_1", vector{m_inputSize1}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "NUM_WG_L_1", ParameterRange(m_inputSize1)); + m_tuner->AddParameter(m_kernel, "NUM_WI_L_1", ParameterRange(m_inputSize1)); + + m_tuner->AddParameter(m_kernel, "INPUT_SIZE_R_1", vector{m_inputSize2}); + m_tuner->AddParameter(m_kernel, "L_CB_SIZE_R_1", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "P_CB_SIZE_R_1", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "NUM_WG_R_1", ParameterRange(m_inputSize2)); + m_tuner->AddParameter(m_kernel, "NUM_WI_R_1", ParameterRange(m_inputSize2)); + + m_tuner->AddParameter(m_kernel, "L_REDUCTION", vector{1}); + m_tuner->AddParameter(m_kernel, "P_WRITE_BACK", vector{0}); + m_tuner->AddParameter(m_kernel, "L_WRITE_BACK", vector{1}); // Add constraints - m_tuner.AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); - m_tuner.AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_R_1"}, UnequalConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); - - m_tuner.AddConstraint(m_kernel, {"L_CB_SIZE_R_1", "INPUT_SIZE_R_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"P_CB_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_R_1", "INPUT_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_R_1", "L_CB_SIZE_R_1", "P_CB_SIZE_R_1"}, DividesDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WI_R_1", "INPUT_SIZE_R_1", "NUM_WG_R_1"}, LessThanOrEqualCeilDivConstraint); - m_tuner.AddConstraint(m_kernel, {"NUM_WG_R_1", "L_CB_SIZE_R_1"}, NoPostInSecondKernelConstraint); + m_tuner->AddConstraint(m_kernel, {"G_CB_RES_DEST_LEVEL", "L_CB_RES_DEST_LEVEL", "P_CB_RES_DEST_LEVEL"}, DescendingConstraint); + m_tuner->AddConstraint(m_kernel, {"OCL_DIM_L_1", "OCL_DIM_R_1"}, UnequalConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_L_1", "INPUT_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_L_1", "INPUT_SIZE_L_1", "L_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "L_CB_SIZE_L_1", "P_CB_SIZE_L_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_L_1", "INPUT_SIZE_L_1", "NUM_WG_L_1"}, LessThanOrEqualCeilDivConstraint); + + m_tuner->AddConstraint(m_kernel, {"L_CB_SIZE_R_1", "INPUT_SIZE_R_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"P_CB_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_R_1", "INPUT_SIZE_R_1", "L_CB_SIZE_R_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_R_1", "L_CB_SIZE_R_1", "P_CB_SIZE_R_1"}, DividesDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WI_R_1", "INPUT_SIZE_R_1", "NUM_WG_R_1"}, LessThanOrEqualCeilDivConstraint); + m_tuner->AddConstraint(m_kernel, {"NUM_WG_R_1", "L_CB_SIZE_R_1"}, NoPostInSecondKernelConstraint); } }; diff --git a/Examples/Bicg/Bicg.cpp b/Examples/Bicg/Bicg.cpp index 9f0e19c5..c32a19ba 100644 --- a/Examples/Bicg/Bicg.cpp +++ b/Examples/Bicg/Bicg.cpp @@ -6,9 +6,9 @@ using namespace std; class Bicg : public ExampleReferenceComputation { protected: - Bicg(shared_ptr config, int defaultProblemSize, string exampleFolderPath, + Bicg(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleReferenceComputation(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName), + ExampleReferenceComputation(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName), // Bicg has O(m × n) complexity. For square matrices where m = n, // we scale with square root of problem size to keep total work proportional m_m(static_cast(sqrt(m_problemSize)) * 1024), @@ -58,13 +58,13 @@ class Bicg : public ExampleReferenceComputation { void InitKernel() override { - m_mId = m_tuner.AddArgumentScalar(m_m); - m_nId = m_tuner.AddArgumentScalar(m_n); - m_AId = m_tuner.AddArgumentVector(m_A, ktt::ArgumentAccessType::ReadWrite); - m_x1Id = m_tuner.AddArgumentVector(m_x1, ktt::ArgumentAccessType::ReadOnly); - m_x2Id = m_tuner.AddArgumentVector(m_x2, ktt::ArgumentAccessType::ReadOnly); - m_y1Id = m_tuner.AddArgumentVector(m_y1, ktt::ArgumentAccessType::ReadWrite); - m_y2Id = m_tuner.AddArgumentVector(m_y2, ktt::ArgumentAccessType::ReadWrite); + m_mId = m_tuner->AddArgumentScalar(m_m); + m_nId = m_tuner->AddArgumentScalar(m_n); + m_AId = m_tuner->AddArgumentVector(m_A, ktt::ArgumentAccessType::ReadWrite); + m_x1Id = m_tuner->AddArgumentVector(m_x1, ktt::ArgumentAccessType::ReadOnly); + m_x2Id = m_tuner->AddArgumentVector(m_x2, ktt::ArgumentAccessType::ReadOnly); + m_y1Id = m_tuner->AddArgumentVector(m_y1, ktt::ArgumentAccessType::ReadWrite); + m_y2Id = m_tuner->AddArgumentVector(m_y2, ktt::ArgumentAccessType::ReadWrite); const ktt::DimensionVector ndRangeDimensions(m_m, m_n / 64); const ktt::DimensionVector workGroupDimensions(32, 4); @@ -72,15 +72,15 @@ class Bicg : public ExampleReferenceComputation { const ktt::DimensionVector referenceNdRangeDimensions2(static_cast(ceil(m_m * 1. / WORK_GROUP_X)), 1); const ktt::DimensionVector referenceWorkGroupDimensions(WORK_GROUP_X, WORK_GROUP_Y); - m_definitionFused = m_tuner.AddKernelDefinitionFromFile("bicgFused", m_kernelFile, ndRangeDimensions, workGroupDimensions); - m_definitionReduction1 = m_tuner.AddKernelDefinitionFromFile("bicgReduction1", m_kernelFile, referenceNdRangeDimensions1, referenceWorkGroupDimensions); - m_definitionReduction2 = m_tuner.AddKernelDefinitionFromFile("bicgReduction2", m_kernelFile, referenceNdRangeDimensions1, referenceWorkGroupDimensions); + m_definitionFused = m_tuner->AddKernelDefinitionFromFile("bicgFused", m_kernelFile, ndRangeDimensions, workGroupDimensions); + m_definitionReduction1 = m_tuner->AddKernelDefinitionFromFile("bicgReduction1", m_kernelFile, referenceNdRangeDimensions1, referenceWorkGroupDimensions); + m_definitionReduction2 = m_tuner->AddKernelDefinitionFromFile("bicgReduction2", m_kernelFile, referenceNdRangeDimensions1, referenceWorkGroupDimensions); - m_kernel = m_tuner.CreateCompositeKernel("BicgPolyBenchAndFused", {m_definitionFused, m_definitionReduction1, m_definitionReduction2}, + m_kernel = m_tuner->CreateCompositeKernel("BicgPolyBenchAndFused", {m_definitionFused, m_definitionReduction1, m_definitionReduction2}, [this](ktt::ComputeInterface& interface) { const vector& parameterValues = interface.GetCurrentConfiguration().GetPairs(); - if (!m_config->useProfiling) + if (!m_useProfiling) { interface.RunKernel(m_definitionFused); } @@ -96,49 +96,49 @@ class Bicg : public ExampleReferenceComputation { } }); - m_tuner.SetArguments(m_definitionFused, {m_AId, m_x1Id, m_y1Id, m_x2Id, m_y2Id, m_mId, m_nId}); - m_tuner.SetArguments(m_definitionReduction1, {m_mId, m_nId, m_y1Id}); - m_tuner.SetArguments(m_definitionReduction2, {m_mId, m_nId, m_y2Id}); + m_tuner->SetArguments(m_definitionFused, {m_AId, m_x1Id, m_y1Id, m_x2Id, m_y2Id, m_mId, m_nId}); + m_tuner->SetArguments(m_definitionReduction1, {m_mId, m_nId, m_y1Id}); + m_tuner->SetArguments(m_definitionReduction2, {m_mId, m_nId, m_y2Id}); } void InitTuningSpace() override { - m_tuner.AddParameter(m_kernel, "FUSED", vector{2}); - m_tuner.AddParameter(m_kernel, "BICG_BATCH", vector{1, 2, 4, 8, 16, 32, 64}); - m_tuner.AddParameter(m_kernel, "USE_SHARED_MATRIX", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "USE_SHARED_VECTOR_1", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "USE_SHARED_VECTOR_2", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "USE_SHARED_REDUCTION_1", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "USE_SHARED_REDUCTION_2", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "ATOMICS", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "UNROLL_BICG_STEP", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "ROWS_PROCESSED", vector{128, 256, 512, 1024}); - m_tuner.AddParameter(m_kernel, "TILE", vector{16, 32, 64}); + m_tuner->AddParameter(m_kernel, "FUSED", vector{2}); + m_tuner->AddParameter(m_kernel, "BICG_BATCH", vector{1, 2, 4, 8, 16, 32, 64}); + m_tuner->AddParameter(m_kernel, "USE_SHARED_MATRIX", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "USE_SHARED_VECTOR_1", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "USE_SHARED_VECTOR_2", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "USE_SHARED_REDUCTION_1", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "USE_SHARED_REDUCTION_2", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "ATOMICS", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "UNROLL_BICG_STEP", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "ROWS_PROCESSED", vector{128, 256, 512, 1024}); + m_tuner->AddParameter(m_kernel, "TILE", vector{16, 32, 64}); auto globalModifierX = [m = m_m](const uint64_t, const vector& v) {return m / v.at(0);}; auto globalModifierY = [n = m_n](const uint64_t, const vector& v) {return n / v.at(0);}; auto localModifierX = [](const uint64_t, const vector& v) {return v.at(0);}; auto localModifierY = [](const uint64_t, const vector& v) {return v.at(0) / v.at(1);}; - m_tuner.AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"TILE"}, globalModifierX); - m_tuner.AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"ROWS_PROCESSED"}, globalModifierY); - m_tuner.AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"TILE"}, localModifierX); - m_tuner.AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"TILE", "BICG_BATCH"}, localModifierY); + m_tuner->AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"TILE"}, globalModifierX); + m_tuner->AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"ROWS_PROCESSED"}, globalModifierY); + m_tuner->AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"TILE"}, localModifierX); + m_tuner->AddThreadModifier(m_kernel, {m_definitionFused}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"TILE", "BICG_BATCH"}, localModifierY); auto fused = [](const vector& v) {return v.at(0) == 2 || ((v.at(0) == 0 || v.at(0) == 1) && v.at(1) == 4 && v.at(2) == 1 && v.at(3) == 1 && v.at(4) == 1 && v.at(5) == 1 && v.at(6) == 1 && v.at(7) == 1 && v.at(8) == 1 && v.at(9) == 512 && v.at(10) == 32); }; - m_tuner.AddConstraint(m_kernel, {"FUSED", "BICG_BATCH", "USE_SHARED_MATRIX", "USE_SHARED_VECTOR_1", "USE_SHARED_VECTOR_2", "USE_SHARED_REDUCTION_1", "USE_SHARED_REDUCTION_2", "ATOMICS", "UNROLL_BICG_STEP", "ROWS_PROCESSED", "TILE"}, fused); + m_tuner->AddConstraint(m_kernel, {"FUSED", "BICG_BATCH", "USE_SHARED_MATRIX", "USE_SHARED_VECTOR_1", "USE_SHARED_VECTOR_2", "USE_SHARED_REDUCTION_1", "USE_SHARED_REDUCTION_2", "ATOMICS", "UNROLL_BICG_STEP", "ROWS_PROCESSED", "TILE"}, fused); auto maxWgSize = [this](const vector& v) {return (v.at(0) * v.at(0) / v.at(1) <= MAX_WORK_GROUP_SIZE) && (v.at(1) <= v.at(0)); }; - m_tuner.AddConstraint(m_kernel, {"TILE", "BICG_BATCH"}, maxWgSize); + m_tuner->AddConstraint(m_kernel, {"TILE", "BICG_BATCH"}, maxWgSize); } void InitReference() override { - m_tuner.SetValidationMethod(ktt::ValidationMethod::SideBySideRelativeComparison, 0.001); - m_tuner.SetValidationRange(m_y1Id, m_n); - m_tuner.SetValidationRange(m_y2Id, m_m); + m_tuner->SetValidationMethod(ktt::ValidationMethod::SideBySideRelativeComparison, 0.001); + m_tuner->SetValidationRange(m_y1Id, m_n); + m_tuner->SetValidationRange(m_y2Id, m_m); - m_tuner.SetReferenceComputation(m_y1Id, [this](void* buffer) + m_tuner->SetReferenceComputation(m_y1Id, [this](void* buffer) { float* y1 = static_cast(buffer); @@ -152,7 +152,7 @@ class Bicg : public ExampleReferenceComputation { } }); - m_tuner.SetReferenceComputation(m_y2Id, [this](void* buffer) + m_tuner->SetReferenceComputation(m_y2Id, [this](void* buffer) { float* y2 = static_cast(buffer); diff --git a/Examples/ClTuneConvolution/ClTuneConvolution.cpp b/Examples/ClTuneConvolution/ClTuneConvolution.cpp index 60f65b85..89571317 100644 --- a/Examples/ClTuneConvolution/ClTuneConvolution.cpp +++ b/Examples/ClTuneConvolution/ClTuneConvolution.cpp @@ -6,9 +6,9 @@ using namespace std; class ClTuneConvolution : public ExampleReferenceKernel { public: - ClTuneConvolution(shared_ptr config, int defaultProblemSize, string exampleFolderPath, + ClTuneConvolution(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName, string defaultRefKernelFileBaseName) : - ExampleReferenceKernel(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName, defaultRefKernelFileBaseName), + ExampleReferenceKernel(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName, defaultRefKernelFileBaseName), kSizeX(4096), kSizeY(4096) { @@ -72,21 +72,21 @@ class ClTuneConvolution : public ExampleReferenceKernel { void InitKernel() override { - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); - kSizeXId = m_tuner.AddArgumentScalar(kSizeX); - kSizeYId = m_tuner.AddArgumentScalar(kSizeY); - matAId = m_tuner.AddArgumentVector(mat_a, ktt::ArgumentAccessType::ReadOnly); - coeffId = m_tuner.AddArgumentVector(coeff, ktt::ArgumentAccessType::ReadOnly); - matBId = m_tuner.AddArgumentVector(mat_b, ktt::ArgumentAccessType::WriteOnly); + m_tuner->SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); + kSizeXId = m_tuner->AddArgumentScalar(kSizeX); + kSizeYId = m_tuner->AddArgumentScalar(kSizeY); + matAId = m_tuner->AddArgumentVector(mat_a, ktt::ArgumentAccessType::ReadOnly); + coeffId = m_tuner->AddArgumentVector(coeff, ktt::ArgumentAccessType::ReadOnly); + matBId = m_tuner->AddArgumentVector(mat_b, ktt::ArgumentAccessType::WriteOnly); const ktt::DimensionVector ndRangeDimensions(kSizeX, kSizeY); const ktt::DimensionVector workGroupDimensions; - m_definition = m_tuner.AddKernelDefinitionFromFile("conv", m_kernelFile, ndRangeDimensions, workGroupDimensions); + m_definition = m_tuner->AddKernelDefinitionFromFile("conv", m_kernelFile, ndRangeDimensions, workGroupDimensions); - m_kernel = m_tuner.CreateSimpleKernel("Convolution", m_definition); + m_kernel = m_tuner->CreateSimpleKernel("Convolution", m_definition); - m_tuner.SetArguments(m_definition, {kSizeXId, kSizeYId, matAId, coeffId, matBId}); + m_tuner->SetArguments(m_definition, {kSizeXId, kSizeYId, matAId, coeffId, matBId}); } void InitTuningSpace() override @@ -94,15 +94,15 @@ class ClTuneConvolution : public ExampleReferenceKernel { vector blockRange = {8, 16, 32, 64}; vector wptRange = {1, 2, 4, 8, 16}; - m_tuner.AddParameter(m_kernel, "TBX", blockRange); - m_tuner.AddParameter(m_kernel, "TBY", blockRange); - m_tuner.AddParameter(m_kernel, "LOCAL", vector{0, 1, 2}); - m_tuner.AddParameter(m_kernel, "WPTX", wptRange); - m_tuner.AddParameter(m_kernel, "WPTY", wptRange); - m_tuner.AddParameter(m_kernel, "VECTOR", vector{1, 2, 4}); - m_tuner.AddParameter(m_kernel, "UNROLL_FACTOR1", vector{0, 1, static_cast(FS)}); - m_tuner.AddParameter(m_kernel, "UNROLL_FACTOR2", vector{0, 1, static_cast(FS)}); - m_tuner.AddParameter(m_kernel, "PADDING", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "TBX", blockRange); + m_tuner->AddParameter(m_kernel, "TBY", blockRange); + m_tuner->AddParameter(m_kernel, "LOCAL", vector{0, 1, 2}); + m_tuner->AddParameter(m_kernel, "WPTX", wptRange); + m_tuner->AddParameter(m_kernel, "WPTY", wptRange); + m_tuner->AddParameter(m_kernel, "VECTOR", vector{1, 2, 4}); + m_tuner->AddParameter(m_kernel, "UNROLL_FACTOR1", vector{0, 1, static_cast(FS)}); + m_tuner->AddParameter(m_kernel, "UNROLL_FACTOR2", vector{0, 1, static_cast(FS)}); + m_tuner->AddParameter(m_kernel, "PADDING", vector{0, 1}); vector integers{8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74}; @@ -111,19 +111,19 @@ class ClTuneConvolution : public ExampleReferenceKernel { // In this case, the workgroup size (TBX by TBY) is extra large (TBX_XL by TBY_XL) because it uses // extra threads to compute the halo threads. How many extra threads are needed is dependend on // the filter size. Here we support a the TBX and TBY size plus up to 10 extra threads. - m_tuner.AddParameter(m_kernel, "TBX_XL", integers); - m_tuner.AddParameter(m_kernel, "TBY_XL", integers); + m_tuner->AddParameter(m_kernel, "TBX_XL", integers); + m_tuner->AddParameter(m_kernel, "TBY_XL", integers); // Add kernel dimension modifiers based on added tuning parameters auto globalModifier = [](const uint64_t size, const vector& v) { return (size * v[0]) / (v[1] * v[2]); }; - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"TBX_XL", "TBX", "WPTX"}, globalModifier); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"TBY_XL", "TBY", "WPTY"}, globalModifier); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"TBX_XL", "TBX", "WPTX"}, globalModifier); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"TBY_XL", "TBY", "WPTY"}, globalModifier); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "TBX_XL", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "TBY_XL", ktt::ModifierAction::Multiply); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "TBX_XL", ktt::ModifierAction::Multiply); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "TBY_XL", ktt::ModifierAction::Multiply); // Add constraints auto haloThreads = [this](const vector& v) { @@ -131,8 +131,8 @@ class ClTuneConvolution : public ExampleReferenceKernel { else { return (v[1] == v[2]); } // Without halo threads }; - m_tuner.AddConstraint(m_kernel, {"LOCAL", "TBX_XL", "TBX", "WPTX"}, haloThreads); - m_tuner.AddConstraint(m_kernel, {"LOCAL", "TBY_XL", "TBY", "WPTY"}, haloThreads); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "TBX_XL", "TBX", "WPTX"}, haloThreads); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "TBY_XL", "TBY", "WPTY"}, haloThreads); // Sets the constrains on the vector size auto vectorConstraint = [this](const vector& v) { @@ -140,16 +140,16 @@ class ClTuneConvolution : public ExampleReferenceKernel { else { return IsMultiple(v[2], v[1]); } }; - m_tuner.AddConstraint(m_kernel, {"LOCAL", "VECTOR", "WPTX"}, vectorConstraint); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "VECTOR", "WPTX"}, vectorConstraint); // Sets padding to zero in case local memory is not used auto paddingConstraint = [](const vector& v) { return (v[1] == 0 || v[0] != 0); }; - m_tuner.AddConstraint(m_kernel, {"LOCAL", "PADDING"}, paddingConstraint); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "PADDING"}, paddingConstraint); // Ensure divisibility auto divConstraint = [](const vector& v) { return v[0] % v[1] == 0; }; - m_tuner.AddConstraint(m_kernel, {"TBX", "WPTX"}, divConstraint); - m_tuner.AddConstraint(m_kernel, {"TBY", "WPTY"}, divConstraint); + m_tuner->AddConstraint(m_kernel, {"TBX", "WPTX"}, divConstraint); + m_tuner->AddConstraint(m_kernel, {"TBY", "WPTY"}, divConstraint); } void InitReference() override diff --git a/Examples/ClTuneGemm/ClTuneGemm.cpp b/Examples/ClTuneGemm/ClTuneGemm.cpp index 5bb175fc..8bab453a 100644 --- a/Examples/ClTuneGemm/ClTuneGemm.cpp +++ b/Examples/ClTuneGemm/ClTuneGemm.cpp @@ -10,9 +10,9 @@ bool IsMultiple(const size_t a, const size_t b) class ClTuneGemm : public ExampleReferenceKernel { protected: - ClTuneGemm(std::shared_ptr config, int defaultProblemSize, string exampleFolderPath, + ClTuneGemm(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName, string defaultRefKernelFileBaseName) : - ExampleReferenceKernel(config, defaultProblemSize, exampleFolderPath, + ExampleReferenceKernel(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName, defaultRefKernelFileBaseName), // GEMM has O(m × n × k) complexity. For square matrices where m = n = k, // we scale with cube root of problem size to keep total work proportional @@ -53,12 +53,12 @@ class ClTuneGemm : public ExampleReferenceKernel { void InitKernel() override { - m_kSizeMId = m_tuner.AddArgumentScalar(m_kSizeM); - m_kSizeNId = m_tuner.AddArgumentScalar(m_kSizeN); - m_kSizeKId = m_tuner.AddArgumentScalar(m_kSizeK); - m_matAId = m_tuner.AddArgumentVector(m_matA, ktt::ArgumentAccessType::ReadOnly); - m_matBId = m_tuner.AddArgumentVector(m_matB, ktt::ArgumentAccessType::ReadOnly); - m_matCId = m_tuner.AddArgumentVector(m_matC, ktt::ArgumentAccessType::WriteOnly); + m_kSizeMId = m_tuner->AddArgumentScalar(m_kSizeM); + m_kSizeNId = m_tuner->AddArgumentScalar(m_kSizeN); + m_kSizeKId = m_tuner->AddArgumentScalar(m_kSizeK); + m_matAId = m_tuner->AddArgumentVector(m_matA, ktt::ArgumentAccessType::ReadOnly); + m_matBId = m_tuner->AddArgumentVector(m_matB, ktt::ArgumentAccessType::ReadOnly); + m_matCId = m_tuner->AddArgumentVector(m_matC, ktt::ArgumentAccessType::WriteOnly); InitKernelDefault("gemm_fast", "Gemm", m_gridDimensions, {m_kSizeMId, m_kSizeNId, m_kSizeKId, m_matAId, m_matBId, m_matCId}); @@ -66,37 +66,37 @@ class ClTuneGemm : public ExampleReferenceKernel { void InitTuningSpace() override { - m_tuner.AddParameter(m_kernel, "MWG", vector{16, 32, 64, 128}); - m_tuner.AddParameter(m_kernel, "NWG", vector{16, 32, 64, 128}); - m_tuner.AddParameter(m_kernel, "KWG", vector{16, 32}); - m_tuner.AddParameter(m_kernel, "MDIMC", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "NDIMC", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "MDIMA", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "NDIMB", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "KWI", vector{2, 8}); + m_tuner->AddParameter(m_kernel, "MWG", vector{16, 32, 64, 128}); + m_tuner->AddParameter(m_kernel, "NWG", vector{16, 32, 64, 128}); + m_tuner->AddParameter(m_kernel, "KWG", vector{16, 32}); + m_tuner->AddParameter(m_kernel, "MDIMC", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "NDIMC", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "MDIMA", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "NDIMB", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "KWI", vector{2, 8}); if (m_computeApi == ktt::ComputeApi::OpenCL) { - m_tuner.AddParameter(m_kernel, "VWM", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "VWN", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "VWM", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "VWN", vector{1, 2, 4, 8}); } else { - m_tuner.AddParameter(m_kernel, "VWM", vector{1, 2, 4}); - m_tuner.AddParameter(m_kernel, "VWN", vector{1, 2, 4}); + m_tuner->AddParameter(m_kernel, "VWM", vector{1, 2, 4}); + m_tuner->AddParameter(m_kernel, "VWN", vector{1, 2, 4}); } - m_tuner.AddParameter(m_kernel, "STRM", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "STRN", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "SA", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "SB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "PRECISION", vector{32}); + m_tuner->AddParameter(m_kernel, "STRM", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "STRN", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "SA", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "SB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "PRECISION", vector{32}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "MWG", ktt::ModifierAction::Divide); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, "NWG", ktt::ModifierAction::Divide); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "MWG", ktt::ModifierAction::Divide); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, "NWG", ktt::ModifierAction::Divide); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "MDIMC", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "NDIMC", ktt::ModifierAction::Multiply); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "MDIMC", ktt::ModifierAction::Multiply); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "NDIMC", ktt::ModifierAction::Multiply); // Add conditions // Sets constraints: Set-up the constraints functions to use. The constraints require a function @@ -108,13 +108,13 @@ class ClTuneGemm : public ExampleReferenceKernel { auto multipleOfXMulY = [](const std::vector& v) {return IsMultiple(v[0], v[1] * v[2]);}; auto multipleOfXMulYDivZ = [](const std::vector& v) {return IsMultiple(v[0], (v[1] * v[2]) / v[3]);}; - m_tuner.AddConstraint(m_kernel, {"KWG", "KWI"}, multipleOfX); - m_tuner.AddConstraint(m_kernel, {"MWG", "MDIMC", "VWM"}, multipleOfXMulY); - m_tuner.AddConstraint(m_kernel, {"NWG", "NDIMC", "VWN"}, multipleOfXMulY); - m_tuner.AddConstraint(m_kernel, {"MWG", "MDIMA", "VWM"}, multipleOfXMulY); - m_tuner.AddConstraint(m_kernel, {"NWG", "NDIMB", "VWN"}, multipleOfXMulY); - m_tuner.AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "MDIMA"}, multipleOfXMulYDivZ); - m_tuner.AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "NDIMB"}, multipleOfXMulYDivZ); + m_tuner->AddConstraint(m_kernel, {"KWG", "KWI"}, multipleOfX); + m_tuner->AddConstraint(m_kernel, {"MWG", "MDIMC", "VWM"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"NWG", "NDIMC", "VWN"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"MWG", "MDIMA", "VWM"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"NWG", "NDIMB", "VWN"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "MDIMA"}, multipleOfXMulYDivZ); + m_tuner->AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "NDIMB"}, multipleOfXMulYDivZ); } void InitReference() override diff --git a/Examples/CliComponent.cpp b/Examples/CliComponent.cpp new file mode 100644 index 00000000..3429ed33 --- /dev/null +++ b/Examples/CliComponent.cpp @@ -0,0 +1,67 @@ +#include "CliComponent.h" +#include +#include +#include +#include + +using namespace std; + +CliOption::CliOption(function &)> callback, const string &trigger, const string &description, + const string &argumentDescriptions, const int argumentCount) + : m_callback(callback), m_trigger(trigger), m_description(description), + m_argumentDescriptions(argumentDescriptions), m_argumentCount(argumentCount) +{ +} + +string CliOption::get_string() const +{ + return m_trigger + " " + m_argumentDescriptions + "\n\t" + m_description; +} + +bool CliOption::TryTrigger(int argc, char **argv, int &i) const { + assert(i < argc); + if (argv[i] != m_trigger) return false; + if (i + m_argumentCount >= argc) + { + cerr << m_trigger << " expects a value to be passed!" << endl; + exit(1); + } + vector arguments; + for (int j = 0; j < m_argumentCount; ++j) { + arguments.push_back(argv[++i]); + } + m_callback(arguments); + return true; +} + +CliComponent::CliComponent() +{ + AddOption({[this](const vector &) { + cout << "Usage: program [options]" << endl << endl; + cout << "Options:" << endl; + for (const auto& option : m_options) { + cout << option.get_string() << endl; + } + exit(0); + }, "--help", "Show this help message and exit."}); +} + +void CliComponent::AddOption(const CliOption &cliOption) { + m_options.push_back(cliOption); +} + +void CliComponent::ProcessInput(int argc, char **argv) { + for (int i = 1; i < argc; ++i) { + bool triggered = false; + for (const auto& option : m_options) { + if (option.TryTrigger(argc, argv, i)) { + triggered = true; + break; + } + } + if (!triggered) { + cerr << argv[i] << " is not a valid option.\n"; + exit(1); + } + } +} \ No newline at end of file diff --git a/Examples/CliComponent.h b/Examples/CliComponent.h new file mode 100644 index 00000000..e5a7c262 --- /dev/null +++ b/Examples/CliComponent.h @@ -0,0 +1,36 @@ +#pragma once + +#include "Ktt.h" +#include +#include +#include +#include +#include +#include + +class CliOption +{ + std::function &)> m_callback; + const std::string m_trigger; + const std::string m_description; + const std::string m_argumentDescriptions; + const int m_argumentCount; + +public: + CliOption(std::function &)> callback, const std::string &trigger, const std::string &description, + const std::string &argumentDescriptions = "", const int argumentCount = 0); + + std::string get_string() const; + + bool TryTrigger(int argc, char **argv, int &i) const; +}; + +class CliComponent { +public: + CliComponent(); + void AddOption(const CliOption &option); + void ProcessInput(int argc, char **argv); + +protected: + std::vector m_options; +}; \ No newline at end of file diff --git a/Examples/CompilerTuningComponent.cpp b/Examples/CompilerTuningComponent.cpp new file mode 100644 index 00000000..a741b5d4 --- /dev/null +++ b/Examples/CompilerTuningComponent.cpp @@ -0,0 +1,55 @@ +#include "CompilerTuningComponent.h" +#include "CliComponent.h" +#include "RunStats.hpp" +#include + +using namespace std; + +CompilerTuningComponent::CompilerTuningComponent(std::shared_ptr &tuner, ktt::KernelId &kernel): + m_tuner(tuner), + m_kernel(kernel), + useSeparateTuning(false) +{} + +void CompilerTuningComponent::AddCompilerParameter(const string &name, const vector &values) +{ + if (useSeparateTuning) m_tuner->AddSeparateCompilerParameter(m_kernel, name, values); + else m_tuner->AddCompilerParameter(m_kernel, name, values); +} + +void CompilerTuningComponent::InitCLIOptions(CliComponent &cli) { + cli.AddOption({[this](const vector &) { + useSeparateTuning = true; + }, "--sepCompTuning", "Enable separate compiler parameter tuning."}); +} + +void CompilerTuningComponent::Run() { + if (!useSeparateTuning) return; + auto bestConfig = m_tuner->GetBestConfiguration(m_kernel); + if (!bestConfig.IsValid()) { + cout << "No valid configuration was found. Skipping separate compiler parameter tuning.\n"; + return; + } + std::cout << "\nTuning compiler options on top of best kernel configuration..." << std::endl; + const auto optResults = m_tuner->TuneOptions(m_kernel, bestConfig); + m_tuner->SaveResults(optResults, "CoulombSumOptionsOutput", ktt::OutputFormat::JSON); + + RunStats stats; + + for (const auto& optResult : optResults) + { + stats.Update(optResult); + } + + stats.Print("Option tuning stats"); +} + +void NoCompilerTuning::AddCompilerParameter(const string &, const vector &) { + assert(false && "Example needs to call UseCompilerTuning() in the constructor to use this feature."); +} + +void NoCompilerTuning::InitCLIOptions(CliComponent &) { +} + +void NoCompilerTuning::Run() { +} diff --git a/Examples/CompilerTuningComponent.h b/Examples/CompilerTuningComponent.h new file mode 100644 index 00000000..50eea225 --- /dev/null +++ b/Examples/CompilerTuningComponent.h @@ -0,0 +1,28 @@ +#include "CliComponent.h" +#include "Tuner.h" +#include + +class CompilerTuningComponent { +public: + CompilerTuningComponent(std::shared_ptr &tuner, ktt::KernelId &kernel); + + // Public virtual despite NVI guidelines because they would just be thin wrappers otherwise. + // Follows "Do not generalize prematurely" and "you aren't gonna need it" + virtual void InitCLIOptions(CliComponent &cli); + virtual void AddCompilerParameter(const std::string &name, const std::vector &values = {}); + virtual void Run(); + +protected: + std::shared_ptr &m_tuner; + ktt::KernelId &m_kernel; + bool useSeparateTuning; +}; + +class NoCompilerTuning : public CompilerTuningComponent { +public: + using CompilerTuningComponent::CompilerTuningComponent; + + void AddCompilerParameter(const std::string &name, const std::vector &values = {}) override; + void InitCLIOptions(CliComponent &cli) override; + void Run() override; +}; \ No newline at end of file diff --git a/Examples/Convolution3d/Convolution3d.cpp b/Examples/Convolution3d/Convolution3d.cpp index 99e7cb48..7b84312b 100644 --- a/Examples/Convolution3d/Convolution3d.cpp +++ b/Examples/Convolution3d/Convolution3d.cpp @@ -6,11 +6,11 @@ class Convolution3d: public ExampleReferenceComputation { protected: Convolution3d( - std::shared_ptr config, + int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName ): - ExampleReferenceComputation(config, defaultProblemSize, + ExampleReferenceComputation(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName ) { @@ -114,12 +114,12 @@ class Convolution3d: public ExampleReferenceComputation const ktt::DimensionVector workGroupDimensions; // Add 3 kernels to the tuner, one of them acts as reference kernel - m_blockedDefinition = m_tuner.AddKernelDefinitionFromFile("conv", m_kernelFile, ndRangeDimensions, + m_blockedDefinition = m_tuner->AddKernelDefinitionFromFile("conv", m_kernelFile, ndRangeDimensions, workGroupDimensions); - m_slidingPlaneDefinition = m_tuner.AddKernelDefinitionFromFile("conv2", m_kernelFile, ndRangeDimensions, + m_slidingPlaneDefinition = m_tuner->AddKernelDefinitionFromFile("conv2", m_kernelFile, ndRangeDimensions, workGroupDimensions); - m_kernel = m_tuner.CreateCompositeKernel("3D Convolution", {m_blockedDefinition, m_slidingPlaneDefinition}, + m_kernel = m_tuner->CreateCompositeKernel("3D Convolution", {m_blockedDefinition, m_slidingPlaneDefinition}, [this](ktt::ComputeInterface& interface) { const vector& parameterValues = interface.GetCurrentConfiguration().GetPairs(); @@ -136,48 +136,48 @@ class Convolution3d: public ExampleReferenceComputation }); // Add all arguments utilized by kernels - const ktt::ArgumentId widthId = m_tuner.AddArgumentScalar(m_width); - const ktt::ArgumentId heightId = m_tuner.AddArgumentScalar(m_height); - const ktt::ArgumentId depthId = m_tuner.AddArgumentScalar(m_depth); - const ktt::ArgumentId srcId = m_tuner.AddArgumentVector(m_src, ktt::ArgumentAccessType::ReadOnly); - const ktt::ArgumentId coeffId = m_tuner.AddArgumentVector(m_coeff, ktt::ArgumentAccessType::ReadOnly); - m_destId = m_tuner.AddArgumentVector(m_dest, ktt::ArgumentAccessType::WriteOnly); + const ktt::ArgumentId widthId = m_tuner->AddArgumentScalar(m_width); + const ktt::ArgumentId heightId = m_tuner->AddArgumentScalar(m_height); + const ktt::ArgumentId depthId = m_tuner->AddArgumentScalar(m_depth); + const ktt::ArgumentId srcId = m_tuner->AddArgumentVector(m_src, ktt::ArgumentAccessType::ReadOnly); + const ktt::ArgumentId coeffId = m_tuner->AddArgumentVector(m_coeff, ktt::ArgumentAccessType::ReadOnly); + m_destId = m_tuner->AddArgumentVector(m_dest, ktt::ArgumentAccessType::WriteOnly); // Set kernel arguments for both tuned kernel and reference kernel - m_tuner.SetArguments(m_blockedDefinition, {widthId, heightId, srcId, coeffId, m_destId}); - m_tuner.SetArguments(m_slidingPlaneDefinition, {widthId, heightId, depthId, srcId, coeffId, m_destId}); + m_tuner->SetArguments(m_blockedDefinition, {widthId, heightId, srcId, coeffId, m_destId}); + m_tuner->SetArguments(m_slidingPlaneDefinition, {widthId, heightId, depthId, srcId, coeffId, m_destId}); } void InitTuningSpace() override { // Add kernel parameters. // 0 - Blocked kernel, 1 - Sliding plane kernel - m_tuner.AddParameter(m_kernel, "ALGORITHM", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "TBX", vector{8, 16, 32, 64}); - m_tuner.AddParameter(m_kernel, "TBY", vector{8, 16, 32, 64}); - m_tuner.AddParameter(m_kernel, "TBZ", vector{1, 2, 4, 8, 16, 32}); - m_tuner.AddParameter(m_kernel, "LOCAL", vector{0, 1, 2}); - m_tuner.AddParameter(m_kernel, "WPTX", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "WPTY", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "WPTZ", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "VECTOR", vector{1, 2, 4}); - m_tuner.AddParameter(m_kernel, "UNROLL_FACTOR", vector{1, static_cast(m_fs)}); - m_tuner.AddParameter(m_kernel, "CONSTANT_COEFF", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "CACHE_WORK_TO_REGS", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "REVERSE_LOOP_ORDER", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "REVERSE_LOOP_ORDER2", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "REVERSE_LOOP_ORDER3", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "PADDING", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "Z_ITERATIONS", vector{4, 8, 16, 32}); + m_tuner->AddParameter(m_kernel, "ALGORITHM", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "TBX", vector{8, 16, 32, 64}); + m_tuner->AddParameter(m_kernel, "TBY", vector{8, 16, 32, 64}); + m_tuner->AddParameter(m_kernel, "TBZ", vector{1, 2, 4, 8, 16, 32}); + m_tuner->AddParameter(m_kernel, "LOCAL", vector{0, 1, 2}); + m_tuner->AddParameter(m_kernel, "WPTX", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "WPTY", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "WPTZ", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "VECTOR", vector{1, 2, 4}); + m_tuner->AddParameter(m_kernel, "UNROLL_FACTOR", vector{1, static_cast(m_fs)}); + m_tuner->AddParameter(m_kernel, "CONSTANT_COEFF", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "CACHE_WORK_TO_REGS", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "REVERSE_LOOP_ORDER", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "REVERSE_LOOP_ORDER2", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "REVERSE_LOOP_ORDER3", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "PADDING", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "Z_ITERATIONS", vector{4, 8, 16, 32}); // Introduces a helper parameter to compute the proper number of threads for the LOCAL == 2 case. // In this case, the workgroup size (TBX by TBY) is extra large (TBX_XL by TBY_XL) because it uses // extra (halo) threads only to load the padding to local memory - they don't compute. vector integers{1, 2, 3, 4, 8, 9, 10, 16, 17, 18, 32, 33, 34, 64, 65, 66}; - m_tuner.AddParameter(m_kernel, "TBX_XL", integers); - m_tuner.AddParameter(m_kernel, "TBY_XL", integers); - m_tuner.AddParameter(m_kernel, "TBZ_XL", integers); + m_tuner->AddParameter(m_kernel, "TBX_XL", integers); + m_tuner->AddParameter(m_kernel, "TBY_XL", integers); + m_tuner->AddParameter(m_kernel, "TBZ_XL", integers); // Modify XY NDRange size for all kernels auto globalModifier = [](const uint64_t size, const vector& v) @@ -185,13 +185,13 @@ class Convolution3d: public ExampleReferenceComputation return (size / (v[0] * v[1])); }; - m_tuner.AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Global, + m_tuner->AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"TBX", "WPTX"}, globalModifier); - m_tuner.AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Global, + m_tuner->AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"TBY", "WPTY"}, globalModifier); // Modify Z NDRange size for Blocked kernel - m_tuner.AddThreadModifier(m_kernel, {m_blockedDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_blockedDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, {"TBZ", "WPTZ"}, globalModifier); // Modify Z NDRange size for Sliding plane kernel @@ -200,15 +200,15 @@ class Convolution3d: public ExampleReferenceComputation return (size / (v[0] * v[1] * v[2])); }; - m_tuner.AddThreadModifier(m_kernel, {m_slidingPlaneDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, + m_tuner->AddThreadModifier(m_kernel, {m_slidingPlaneDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, {"TBZ", "WPTZ", "Z_ITERATIONS"}, globalModifierZ); // Modify workgroup size for all kernels - m_tuner.AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Local, + m_tuner->AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "TBX_XL", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Local, + m_tuner->AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "TBY_XL", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Local, + m_tuner->AddThreadModifier(m_kernel, {m_blockedDefinition, m_slidingPlaneDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::Z, "TBZ_XL", ktt::ModifierAction::Multiply); // For LOCAL == 2, extend block size by halo threads @@ -224,13 +224,13 @@ class Convolution3d: public ExampleReferenceComputation } }; - m_tuner.AddConstraint(m_kernel, {"LOCAL", "TBX_XL", "TBX", "WPTX"}, HaloThreads); - m_tuner.AddConstraint(m_kernel, {"LOCAL", "TBY_XL", "TBY", "WPTY"}, HaloThreads); - m_tuner.AddConstraint(m_kernel, {"LOCAL", "TBZ_XL", "TBZ", "WPTZ"}, HaloThreads); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "TBX_XL", "TBX", "WPTX"}, HaloThreads); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "TBY_XL", "TBY", "WPTY"}, HaloThreads); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "TBZ_XL", "TBZ", "WPTZ"}, HaloThreads); // Sets padding to zero in case local memory is not used auto padding = [](const vector& v) { return (v[0] != 0 || v[1] == 0); }; - m_tuner.AddConstraint(m_kernel, {"LOCAL", "PADDING"}, padding); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "PADDING"}, padding); // GPUs have max. workgroup size auto maxWgSize = [this](const vector& v) @@ -238,7 +238,7 @@ class Convolution3d: public ExampleReferenceComputation return v[0] * v[1] * v[2] <= m_maxWorkGroupSize; }; - m_tuner.AddConstraint(m_kernel, {"TBX_XL", "TBY_XL", "TBZ_XL"}, maxWgSize); + m_tuner->AddConstraint(m_kernel, {"TBX_XL", "TBY_XL", "TBZ_XL"}, maxWgSize); // GPUs have max. local memory size auto maxLocalMemSize = [this](const vector& v) @@ -249,11 +249,11 @@ class Convolution3d: public ExampleReferenceComputation * sizeof(float) <= m_maxLocalMemorySize; }; - m_tuner.AddConstraint(m_kernel, {"ALGORITHM", "LOCAL", "PADDING", "TBX_XL", "WPTX", "TBY_XL", "WPTY", "TBZ_XL", "WPTZ"}, + m_tuner->AddConstraint(m_kernel, {"ALGORITHM", "LOCAL", "PADDING", "TBX_XL", "WPTX", "TBY_XL", "WPTY", "TBZ_XL", "WPTZ"}, maxLocalMemSize); auto reverseCacheLoopsOrder = [](const vector& v) { return v[0] == 1 || v[1] == 0; }; - m_tuner.AddConstraint(m_kernel, {"CACHE_WORK_TO_REGS", "REVERSE_LOOP_ORDER3"}, reverseCacheLoopsOrder); + m_tuner->AddConstraint(m_kernel, {"CACHE_WORK_TO_REGS", "REVERSE_LOOP_ORDER3"}, reverseCacheLoopsOrder); // Sets the constrains on the vector size auto vectorConstraint = [this](const vector& v) @@ -268,7 +268,7 @@ class Convolution3d: public ExampleReferenceComputation } }; - m_tuner.AddConstraint(m_kernel, {"LOCAL", "VECTOR", "WPTX"}, vectorConstraint); + m_tuner->AddConstraint(m_kernel, {"LOCAL", "VECTOR", "WPTX"}, vectorConstraint); auto algorithm = [](const vector& v) { @@ -284,17 +284,17 @@ class Convolution3d: public ExampleReferenceComputation } }; - m_tuner.AddConstraint(m_kernel, {"ALGORITHM", "TBX", "TBY", "TBZ", "WPTX", "WPTY", "WPTZ", "LOCAL", "VECTOR", "UNROLL_FACTOR", + m_tuner->AddConstraint(m_kernel, {"ALGORITHM", "TBX", "TBY", "TBZ", "WPTX", "WPTY", "WPTZ", "LOCAL", "VECTOR", "UNROLL_FACTOR", "CONSTANT_COEFF", "CACHE_WORK_TO_REGS", "REVERSE_LOOP_ORDER", "REVERSE_LOOP_ORDER2", "REVERSE_LOOP_ORDER3"}, algorithm); auto slidingPlane = [](const vector& v) { return v[0] == 1 || v[1] == 16; }; - m_tuner.AddConstraint(m_kernel, {"ALGORITHM", "Z_ITERATIONS"}, slidingPlane); + m_tuner->AddConstraint(m_kernel, {"ALGORITHM", "Z_ITERATIONS"}, slidingPlane); } void InitReference() override { - m_tuner.SetValidationMethod(ktt::ValidationMethod::SideBySideComparison, 0.001f); - m_tuner.SetReferenceComputation(m_destId, [this](void* buffer) + m_tuner->SetValidationMethod(ktt::ValidationMethod::SideBySideComparison, 0.001f); + m_tuner->SetReferenceComputation(m_destId, [this](void* buffer) { float* output = static_cast(buffer); diff --git a/Examples/CoulombSum2d/CoulombSum2d.cpp b/Examples/CoulombSum2d/CoulombSum2d.cpp index 4f7ea90c..a9344fb9 100644 --- a/Examples/CoulombSum2d/CoulombSum2d.cpp +++ b/Examples/CoulombSum2d/CoulombSum2d.cpp @@ -5,9 +5,9 @@ using namespace std; class CoulombSum2d : public ExampleReferenceKernel { protected: - CoulombSum2d(std::shared_ptr config, int defaultProblemSize, string exampleFolderPath, + CoulombSum2d(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName, string defaultRefKernelFileBaseName) : - ExampleReferenceKernel(config, defaultProblemSize, exampleFolderPath, + ExampleReferenceKernel(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName, defaultRefKernelFileBaseName), // Since CoulombSum2d has O(n²) complexity (gridPoints × atoms), scale grid dimensions // with the fourth root of problem size to keep total work proportional @@ -72,14 +72,14 @@ class CoulombSum2d : public ExampleReferenceKernel { void InitKernel() override { // Add all kernel arguments - m_atomInfoId = m_tuner.AddArgumentVector(m_atomInfo, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoXId = m_tuner.AddArgumentVector(m_atomInfoX, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoYId = m_tuner.AddArgumentVector(m_atomInfoY, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoZId = m_tuner.AddArgumentVector(m_atomInfoZ, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoWId = m_tuner.AddArgumentVector(m_atomInfoW, ktt::ArgumentAccessType::ReadOnly); - m_numberOfAtomsId = m_tuner.AddArgumentScalar(m_numberOfAtoms); - m_gridSpacingId = m_tuner.AddArgumentScalar(m_gridSpacing); - m_energyGridId = m_tuner.AddArgumentVector(m_energyGrid, ktt::ArgumentAccessType::ReadWrite); + m_atomInfoId = m_tuner->AddArgumentVector(m_atomInfo, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoXId = m_tuner->AddArgumentVector(m_atomInfoX, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoYId = m_tuner->AddArgumentVector(m_atomInfoY, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoZId = m_tuner->AddArgumentVector(m_atomInfoZ, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoWId = m_tuner->AddArgumentVector(m_atomInfoW, ktt::ArgumentAccessType::ReadOnly); + m_numberOfAtomsId = m_tuner->AddArgumentScalar(m_numberOfAtoms); + m_gridSpacingId = m_tuner->AddArgumentScalar(m_gridSpacing); + m_energyGridId = m_tuner->AddArgumentVector(m_energyGrid, ktt::ArgumentAccessType::ReadWrite); // Configure main kernel InitKernelDefault("directCoulombSum", "CoulombSum", m_ndRangeDimensions, @@ -91,31 +91,31 @@ class CoulombSum2d : public ExampleReferenceKernel { { UseFastMath(); - m_tuner.AddParameter(m_kernel, "INNER_UNROLL_FACTOR", vector{0, 1, 2, 4, 8, 16, 32}); - m_tuner.AddParameter(m_kernel, "USE_CONSTANT_MEMORY", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "VECTOR_TYPE", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "USE_SOA", vector{0, 1, 2}); + m_tuner->AddParameter(m_kernel, "INNER_UNROLL_FACTOR", vector{0, 1, 2, 4, 8, 16, 32}); + m_tuner->AddParameter(m_kernel, "USE_CONSTANT_MEMORY", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "VECTOR_TYPE", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "USE_SOA", vector{0, 1, 2}); // Using vectorized SoA only makes sense when vectors are longer than 1. auto vectorizedSoA = [](const vector& vector) {return vector[0] > 1 || vector[1] != 2;}; - m_tuner.AddConstraint(m_kernel, {"VECTOR_TYPE", "USE_SOA"}, vectorizedSoA); + m_tuner->AddConstraint(m_kernel, {"VECTOR_TYPE", "USE_SOA"}, vectorizedSoA); // Divide NDRange in dimension x by OUTER_UNROLL_FACTOR. - m_tuner.AddParameter(m_kernel, "OUTER_UNROLL_FACTOR", vector{1, 2, 4, 8}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "OUTER_UNROLL_FACTOR", + m_tuner->AddParameter(m_kernel, "OUTER_UNROLL_FACTOR", vector{1, 2, 4, 8}); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "OUTER_UNROLL_FACTOR", ktt::ModifierAction::Divide); // Multiply work-group size in dimensions x and y by the following parameters (effectively setting work-group size to their values). - m_tuner.AddParameter(m_kernel, "WORK_GROUP_SIZE_X", vector{4, 8, 16, 32}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", + m_tuner->AddParameter(m_kernel, "WORK_GROUP_SIZE_X", vector{4, 8, 16, 32}); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", ktt::ModifierAction::Divide); - m_tuner.AddParameter(m_kernel, "WORK_GROUP_SIZE_Y", vector{1, 2, 4, 8, 16, 32}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", + m_tuner->AddParameter(m_kernel, "WORK_GROUP_SIZE_Y", vector{1, 2, 4, 8, 16, 32}); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", ktt::ModifierAction::Divide); } diff --git a/Examples/CoulombSum3d/CoulombSum3d.cpp b/Examples/CoulombSum3d/CoulombSum3d.cpp index 96f2d3c0..527aa6fa 100644 --- a/Examples/CoulombSum3d/CoulombSum3d.cpp +++ b/Examples/CoulombSum3d/CoulombSum3d.cpp @@ -7,105 +7,21 @@ using namespace std; -struct CoulombSum3dConfiguration : ExampleConfiguration -{ - uint64_t gridWidth = 128; - uint64_t gridHeight = 128; - uint64_t gridDepth = 128; - uint64_t numberOfAtoms = 256; - float gridSpacing = 0.5f; - bool sepCompTuning = false; -}; - -void SetUpCoulombSum3dOptions(vector &options, CoulombSum3dConfiguration &config) -{ - options.emplace_back([&config](const vector &args) { - config.gridWidth = stoul(args[0]); - }, "--gridWidth", "Set the grid width (expects int), scaled proportionally with" - " cbrt of problemSize", "", 1); - options.emplace_back([&config](const vector &args) { - config.gridHeight = stoul(args[0]); - }, "--gridHeight", "Set the grid height (expects int), scaled proportionally with" - " cbrt of problemSize", "", 1); - options.emplace_back([&config](const vector &args) { - config.gridDepth = stoul(args[0]); - }, "--gridDepth", "Set the grid depth (expects int), scaled proportionally with" - " cbrt of problemSize", "", 1); - options.emplace_back([&config](const vector &args) { - config.numberOfAtoms = stoul(args[0]); - }, "--atoms", "Set the number of atoms (expects int)", "", 1); - options.emplace_back([&config](const vector &args) { - config.gridSpacing = stof(args[0]); - }, "--spacing", "Set the grid spacing (expects float)", "", 1); - options.emplace_back([&config](const vector &) { - config.sepCompTuning = true; - }, "--sepCompTuning", "Enable separate compiler parameter tuning."); -} - -CoulombSum3dConfiguration CoulombSum3dProcessInput(int argc, char **argv) -{ - CoulombSum3dConfiguration config; - vector options; - SetUpCommonOptions(options, &config); - SetUpCoulombSum3dOptions(options, config); - - IterateArguments(argc, argv, options); - - return config; -} - class CoulombSum3d : public ExampleReferenceComputation { protected: - CoulombSum3d(shared_ptr config, int defaultProblemSize, string exampleFolderPath, + CoulombSum3d(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleReferenceComputation(config, defaultProblemSize, exampleFolderPath, + ExampleReferenceComputation(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName), - m_gridWidth(config->gridWidth * static_cast(cbrt(m_problemSize))), - m_gridHeight(config->gridHeight * static_cast(cbrt(m_problemSize))), - m_gridDepth(config->gridDepth * static_cast(cbrt(m_problemSize))), + m_gridWidth(128 * static_cast(cbrt(m_problemSize))), + m_gridHeight(128 * static_cast(cbrt(m_problemSize))), + m_gridDepth(128 * static_cast(cbrt(m_problemSize))), m_ndRangeDimensions(m_gridWidth, m_gridHeight, m_gridDepth), m_workGroupDimensions{1, 1, 1}, - m_numberOfAtoms(config->numberOfAtoms), - m_gridSpacing(config->gridSpacing), - m_sepCompTuning(config->sepCompTuning) + m_numberOfAtoms(256), + m_gridSpacing(0.5f) { - } - -public: - static unique_ptr Create( - int argc, char** argv, - int defaultProblemSize, - string exampleFolderPath, - string defaultKernelFileBaseName - ) { - auto config = make_shared(CoulombSum3dProcessInput(argc, argv)); - unique_ptr ex(new CoulombSum3d(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName)); - ex->PostInitialize(); - return ex; - } - - void Run() - { - ExampleBase::Run(); - - if (!m_sepCompTuning) return; - auto bestConfig = m_tuner.GetBestConfiguration(m_kernel); - if (!bestConfig.IsValid()) { - cout << "No valid configuration was found. Skippin separate compiler parameter tuning.\n"; - return; - } - std::cout << "\nTuning compiler options on top of best kernel configuration..." << std::endl; - const auto optResults = m_tuner.TuneOptions(m_kernel, bestConfig); - m_tuner.SaveResults(optResults, "CoulombSumOptionsOutput", ktt::OutputFormat::JSON); - - RunStats stats; - - for (const auto& optResult : optResults) - { - stats.Update(optResult); - } - - PrintRunStats("Option tuning stats", stats); + UseCompilerTuning(); } protected: @@ -120,10 +36,8 @@ class CoulombSum3d : public ExampleReferenceComputation { const ktt::DimensionVector m_workGroupDimensions; const ktt::DimensionVector m_referenceWorkGroupDimensions{16, 16, 16}; - const int m_numberOfAtoms; - const float m_gridSpacing; - - const bool m_sepCompTuning; + int m_numberOfAtoms; + float m_gridSpacing; vector m_atomInfo; vector m_atomInfoX; @@ -142,10 +56,27 @@ class CoulombSum3d : public ExampleReferenceComputation { ktt::ArgumentId m_gridDimId; ktt::ArgumentId m_energyGridId; - void AddCompilerParameter(const string &name, const vector &values = {}) + void InitCLI() override { - if (m_sepCompTuning) m_tuner.AddSeparateCompilerParameter(m_kernel, name, values); - else m_tuner.AddCompilerParameter(m_kernel, name, values); + ExampleBase::InitCLI(); + m_cli.AddOption({[this](const vector &args) { + m_gridWidth = stoul(args[0]); + }, "--gridWidth", "Set the grid width (expects int), scaled proportionally with" + " cbrt of problemSize", "", 1}); + m_cli.AddOption({[this](const vector &args) { + m_gridHeight = stoul(args[0]); + }, "--gridHeight", "Set the grid height (expects int), scaled proportionally with" + " cbrt of problemSize", "", 1}); + m_cli.AddOption({[this](const vector &args) { + m_gridDepth = stoul(args[0]); + }, "--gridDepth", "Set the grid depth (expects int), scaled proportionally with" + " cbrt of problemSize", "", 1}); + m_cli.AddOption({[this](const vector &args) { + m_numberOfAtoms = stoul(args[0]); + }, "--atoms", "Set the number of atoms (expects int)", "", 1}); + m_cli.AddOption({[this](const vector &args) { + m_gridSpacing = stof(args[0]); + }, "--spacing", "Set the grid spacing (expects float)", "", 1}); } void InitData() override @@ -174,15 +105,15 @@ class CoulombSum3d : public ExampleReferenceComputation { void InitKernel() override { // Add all kernel arguments - m_atomInfoId = m_tuner.AddArgumentVector(m_atomInfo, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoXId = m_tuner.AddArgumentVector(m_atomInfoX, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoYId = m_tuner.AddArgumentVector(m_atomInfoY, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoZId = m_tuner.AddArgumentVector(m_atomInfoZ, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoWId = m_tuner.AddArgumentVector(m_atomInfoW, ktt::ArgumentAccessType::ReadOnly); - m_numberOfAtomsId = m_tuner.AddArgumentScalar(m_numberOfAtoms); - m_gridSpacingId = m_tuner.AddArgumentScalar(m_gridSpacing); - m_gridDimId = m_tuner.AddArgumentScalar(static_cast(m_gridWidth)); - m_energyGridId = m_tuner.AddArgumentVector(m_energyGrid, ktt::ArgumentAccessType::WriteOnly); + m_atomInfoId = m_tuner->AddArgumentVector(m_atomInfo, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoXId = m_tuner->AddArgumentVector(m_atomInfoX, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoYId = m_tuner->AddArgumentVector(m_atomInfoY, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoZId = m_tuner->AddArgumentVector(m_atomInfoZ, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoWId = m_tuner->AddArgumentVector(m_atomInfoW, ktt::ArgumentAccessType::ReadOnly); + m_numberOfAtomsId = m_tuner->AddArgumentScalar(m_numberOfAtoms); + m_gridSpacingId = m_tuner->AddArgumentScalar(m_gridSpacing); + m_gridDimId = m_tuner->AddArgumentScalar(static_cast(m_gridWidth)); + m_energyGridId = m_tuner->AddArgumentVector(m_energyGrid, ktt::ArgumentAccessType::WriteOnly); // Configure main kernel InitKernelDefault("directCoulombSum", "CoulombSum", m_ndRangeDimensions, @@ -197,117 +128,114 @@ class CoulombSum3d : public ExampleReferenceComputation { if (m_computeApi == ktt::ComputeApi::OpenCL || m_computeApi == ktt::ComputeApi::CUDA) { - m_tuner.AddParameter(m_kernel, "WORK_GROUP_SIZE_X", vector{16, 32}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", + m_tuner->AddParameter(m_kernel, "WORK_GROUP_SIZE_X", vector{16, 32}); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, "WORK_GROUP_SIZE_X", ktt::ModifierAction::DivideCeil); - m_tuner.AddParameter(m_kernel, "WORK_GROUP_SIZE_Y", vector{1, 2, 4, 8}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", + m_tuner->AddParameter(m_kernel, "WORK_GROUP_SIZE_Y", vector{1, 2, 4, 8}); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", ktt::ModifierAction::Multiply); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, "WORK_GROUP_SIZE_Y", ktt::ModifierAction::DivideCeil); - m_tuner.AddParameter(m_kernel, "WORK_GROUP_SIZE_Z", vector{1}); - m_tuner.AddParameter(m_kernel, "Z_ITERATIONS", vector{1, 2, 4, 8, 16, 32}); - m_tuner.AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, "Z_ITERATIONS", + m_tuner->AddParameter(m_kernel, "WORK_GROUP_SIZE_Z", vector{1}); + m_tuner->AddParameter(m_kernel, "Z_ITERATIONS", vector{1, 2, 4, 8, 16, 32}); + m_tuner->AddThreadModifier(m_kernel, {m_definition}, ktt::ModifierType::Global, ktt::ModifierDimension::Z, "Z_ITERATIONS", ktt::ModifierAction::DivideCeil); - m_tuner.AddParameter(m_kernel, "INNER_UNROLL_FACTOR", vector{0, 1, 2, 4, 8, 16, 32}); + m_tuner->AddParameter(m_kernel, "INNER_UNROLL_FACTOR", vector{0, 1, 2, 4, 8, 16, 32}); auto lt = [](const vector& vector) {return vector.at(0) < vector.at(1);}; - m_tuner.AddConstraint(m_kernel, {"INNER_UNROLL_FACTOR", "Z_ITERATIONS"}, lt); + m_tuner->AddConstraint(m_kernel, {"INNER_UNROLL_FACTOR", "Z_ITERATIONS"}, lt); auto par = [](const vector& vector) {return vector.at(0) * vector.at(1) >= 64;}; - m_tuner.AddConstraint(m_kernel, {"WORK_GROUP_SIZE_X", "WORK_GROUP_SIZE_Y"}, par); + m_tuner->AddConstraint(m_kernel, {"WORK_GROUP_SIZE_X", "WORK_GROUP_SIZE_Y"}, par); if (m_computeApi == ktt::ComputeApi::OpenCL) { - m_tuner.AddParameter(m_kernel, "USE_CONSTANT_MEMORY", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "USE_SOA", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "VECTOR_SIZE", vector{1, 2 , 4, 8, 16}); + m_tuner->AddParameter(m_kernel, "USE_CONSTANT_MEMORY", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "USE_SOA", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "VECTOR_SIZE", vector{1, 2 , 4, 8, 16}); auto vec = [](const vector& vector) {return vector.at(0) || vector.at(1) == 1; }; - m_tuner.AddConstraint(m_kernel, {"USE_SOA", "VECTOR_SIZE"}, vec); + m_tuner->AddConstraint(m_kernel, {"USE_SOA", "VECTOR_SIZE"}, vec); // Individual math optimization flags replacing the global -cl-fast-relaxed-math. // -cl-fast-relaxed-math ≈ -cl-finite-math-only + -cl-unsafe-math-optimizations, // where -cl-unsafe-math-optimizations implies -cl-mad-enable, -cl-no-signed-zeros, // and -cl-denorms-are-zero. Tuning them individually reveals which subset is sufficient. - AddCompilerParameter("-cl-mad-enable"); - AddCompilerParameter("-cl-no-signed-zeros"); - AddCompilerParameter("-cl-finite-math-only"); - AddCompilerParameter("-cl-denorms-are-zero"); + m_compilerTuning->AddCompilerParameter("-cl-mad-enable"); + m_compilerTuning->AddCompilerParameter("-cl-no-signed-zeros"); + m_compilerTuning->AddCompilerParameter("-cl-finite-math-only"); + m_compilerTuning->AddCompilerParameter("-cl-denorms-are-zero"); } else // CUDA { - m_tuner.AddParameter(m_kernel, "USE_CONSTANT_MEMORY", vector{0}); - m_tuner.AddParameter(m_kernel, "USE_SOA", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "VECTOR_SIZE", vector{1}); + m_tuner->AddParameter(m_kernel, "USE_CONSTANT_MEMORY", vector{0}); + m_tuner->AddParameter(m_kernel, "USE_SOA", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "VECTOR_SIZE", vector{1}); // Register count limit: trades register file pressure for higher occupancy. // Relevant here because energyValue[Z_ITERATIONS] can hold up to 32 floats, // creating high register pressure at large Z_ITERATIONS values. // 0 = unlimited (compiler decides), others force a ceiling. - AddCompilerParameter("--maxrregcount ", {"0", "32", "40", "48", "64"}); + m_compilerTuning->AddCompilerParameter("--maxrregcount ", {"0", "32", "40", "48", "64"}); } } else // CPP { - m_tuner.AddParameter(m_kernel, "OMP_COLLAPSE", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "OMP_SCHEDULING", vector{0, 1, 2}); - m_tuner.AddParameter(m_kernel, "OMP_SCHED_CHUNK", vector{2, 4, 8, 16, 32, 64, 128}); - m_tuner.AddParameter(m_kernel, "TILE", vector{0, 8, 16, 32, 64}); + m_tuner->AddParameter(m_kernel, "OMP_COLLAPSE", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "OMP_SCHEDULING", vector{0, 1, 2}); + m_tuner->AddParameter(m_kernel, "OMP_SCHED_CHUNK", vector{2, 4, 8, 16, 32, 64, 128}); + m_tuner->AddParameter(m_kernel, "TILE", vector{0, 8, 16, 32, 64}); - AddCompilerParameter("-ffast-math"); - AddCompilerParameter("-O", {"1", "2", "3"}); - AddCompilerParameter("-funroll-loops"); + m_compilerTuning->AddCompilerParameter("-ffast-math"); + m_compilerTuning->AddCompilerParameter("-O", {"1", "2", "3"}); + m_compilerTuning->AddCompilerParameter("-funroll-loops"); // Auto-vectorization of the inner atom loop (SIMD via AVX2/AVX512 enabled by -march=native). - AddCompilerParameter("-ftree-vectorize"); + m_compilerTuning->AddCompilerParameter("-ftree-vectorize"); // Software prefetching of the atom SoA arrays in the inner loop (GCC-specific). - AddCompilerParameter("-fprefetch-loop-arrays"); + m_compilerTuning->AddCompilerParameter("-fprefetch-loop-arrays"); // Allows the compiler to skip errno updates in math functions (lighter than -ffast-math, // enables vectorization of sqrtf without full unsafe-math semantics). - AddCompilerParameter("-fno-math-errno"); + m_compilerTuning->AddCompilerParameter("-fno-math-errno"); auto schedchunk = [](const vector& vector) {return vector.at(0) == 2 || vector.at(1) == 2; }; - m_tuner.AddConstraint(m_kernel, {"OMP_SCHEDULING", "OMP_SCHED_CHUNK"}, schedchunk); + m_tuner->AddConstraint(m_kernel, {"OMP_SCHEDULING", "OMP_SCHED_CHUNK"}, schedchunk); } } void InitReference() override { - if (!m_config->rapidTest) - { - //TODO: this is temporary hack, there should be composition of zeroizing and Coulomb kernel, - // otherwise, multiple profiling runs corrupt results - m_tuner.SetReferenceComputation(m_energyGridId, [this](void* buffer) { - float* grid = static_cast(buffer); - - #pragma omp parallel for - for (size_t z = 0; z < m_gridDepth; z++) - for (size_t y = 0; y < m_gridHeight; y++) - for (size_t x = 0; x < m_gridWidth; x++) - { - float e = 0.0f; - float gx = static_cast(x) * m_gridSpacing; - float gy = static_cast(y) * m_gridSpacing; - float gz = static_cast(z) * m_gridSpacing; - for (int a = 0; a < m_numberOfAtoms; a++) - e += m_atomInfoW[a] * (1.0f / sqrtf((m_atomInfoX[a] - gx) * (m_atomInfoX[a] - gx) + - (m_atomInfoY[a] - gy) * (m_atomInfoY[a] - gy) + - (m_atomInfoZ[a] - gz) * (m_atomInfoZ[a] - gz))); - grid[z * m_gridWidth * m_gridHeight + y * m_gridWidth + x] = e; - } - }); - m_tuner.SetValidationMethod(ktt::ValidationMethod::SideBySideComparison, 0.01); - } + //TODO: this is temporary hack, there should be composition of zeroizing and Coulomb kernel, + // otherwise, multiple profiling runs corrupt results + m_tuner->SetReferenceComputation(m_energyGridId, [this](void* buffer) { + float* grid = static_cast(buffer); + + #pragma omp parallel for + for (size_t z = 0; z < m_gridDepth; z++) + for (size_t y = 0; y < m_gridHeight; y++) + for (size_t x = 0; x < m_gridWidth; x++) + { + float e = 0.0f; + float gx = static_cast(x) * m_gridSpacing; + float gy = static_cast(y) * m_gridSpacing; + float gz = static_cast(z) * m_gridSpacing; + for (int a = 0; a < m_numberOfAtoms; a++) + e += m_atomInfoW[a] * (1.0f / sqrtf((m_atomInfoX[a] - gx) * (m_atomInfoX[a] - gx) + + (m_atomInfoY[a] - gy) * (m_atomInfoY[a] - gy) + + (m_atomInfoZ[a] - gz) * (m_atomInfoZ[a] - gz))); + grid[z * m_gridWidth * m_gridHeight + y * m_gridWidth + x] = e; + } + }); + m_tuner->SetValidationMethod(ktt::ValidationMethod::SideBySideComparison, 0.01); } }; int main(int argc, char **argv) { - unique_ptr coulombSum3d = CoulombSum3d::Create(argc, argv, 8, "Examples/CoulombSum3d", "CoulombSum3d"); + unique_ptr coulombSum3d = CoulombSum3d::Create(argc, argv, 8, "Examples/CoulombSum3d", "CoulombSum3d"); coulombSum3d->Run(); return 0; diff --git a/Examples/Covariance/Covariance.cpp b/Examples/Covariance/Covariance.cpp index 743d34fe..089eb745 100644 --- a/Examples/Covariance/Covariance.cpp +++ b/Examples/Covariance/Covariance.cpp @@ -7,9 +7,9 @@ bool IsMultiple(size_t a, size_t b) { return a % b == 0; }; class Covariance : public ExampleReferenceComputation { protected: - Covariance(shared_ptr config, int defaultProblemSize, string exampleFolderPath, + Covariance(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleReferenceComputation(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName), + ExampleReferenceComputation(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName), // Covariance has O(n × m²) complexity. For square matrices where n ≈ m, // we scale with cube root of problem size to keep total work proportional m_n(static_cast(sqrt(m_problemSize)) * 1024), @@ -17,7 +17,7 @@ class Covariance : public ExampleReferenceComputation { { m_refKernelFile = GetKernelFilePath(exampleFolderPath, "CovarianceReference"); m_gemmFile = GetKernelFilePath(exampleFolderPath, "Gemm"); - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); + m_tuner->SetGlobalSizeType(ktt::GlobalSizeType::OpenCL); } friend ExampleBase; @@ -63,29 +63,29 @@ class Covariance : public ExampleReferenceComputation { void InitKernel() override { const float floatN = static_cast(m_n); - mMId = m_tuner.AddArgumentScalar(m_m); - mNId = m_tuner.AddArgumentScalar(m_n); - mFloatNId = m_tuner.AddArgumentScalar(floatN); - mDataId = m_tuner.AddArgumentVector(m_data, ktt::ArgumentAccessType::ReadWrite); - mSymmatId = m_tuner.AddArgumentVector(m_symmat, ktt::ArgumentAccessType::ReadWrite); - mMeanId = m_tuner.AddArgumentVector(m_mean, ktt::ArgumentAccessType::ReadWrite); + mMId = m_tuner->AddArgumentScalar(m_m); + mNId = m_tuner->AddArgumentScalar(m_n); + mFloatNId = m_tuner->AddArgumentScalar(floatN); + mDataId = m_tuner->AddArgumentVector(m_data, ktt::ArgumentAccessType::ReadWrite); + mSymmatId = m_tuner->AddArgumentVector(m_symmat, ktt::ArgumentAccessType::ReadWrite); + mMeanId = m_tuner->AddArgumentVector(m_mean, ktt::ArgumentAccessType::ReadWrite); const ktt::DimensionVector ndRangeDim1D(m_m, 1); const ktt::DimensionVector workGroupDim1D(256, 1); const ktt::DimensionVector ndRangeDim2D(m_m, m_m); const ktt::DimensionVector workGroupDim2D(32, 8); - m_refMeanDefinition = m_tuner.AddKernelDefinitionFromFile("mean_kernel_reference", m_refKernelFile, ndRangeDim1D, workGroupDim1D); - m_refReduceDefinition = m_tuner.AddKernelDefinitionFromFile("reduce_kernel_reference", m_refKernelFile, ndRangeDim2D, workGroupDim2D); - m_refCovarDefinition = m_tuner.AddKernelDefinitionFromFile("covar_kernel_reference", m_refKernelFile, ndRangeDim1D, workGroupDim1D); + m_refMeanDefinition = m_tuner->AddKernelDefinitionFromFile("mean_kernel_reference", m_refKernelFile, ndRangeDim1D, workGroupDim1D); + m_refReduceDefinition = m_tuner->AddKernelDefinitionFromFile("reduce_kernel_reference", m_refKernelFile, ndRangeDim2D, workGroupDim2D); + m_refCovarDefinition = m_tuner->AddKernelDefinitionFromFile("covar_kernel_reference", m_refKernelFile, ndRangeDim1D, workGroupDim1D); - m_meanDefinition = m_tuner.AddKernelDefinitionFromFile("mean_kernel", m_kernelFile, ndRangeDim1D, workGroupDim1D); - m_reduceDefinition = m_tuner.AddKernelDefinitionFromFile("reduce_kernel", m_kernelFile, ndRangeDim2D, workGroupDim2D); - m_covarDefinition = m_tuner.AddKernelDefinitionFromFile("covar_kernel", m_kernelFile, ndRangeDim1D, workGroupDim1D); - m_gemmDefinition = m_tuner.AddKernelDefinitionFromFile("gemm_fast", m_gemmFile, ndRangeDim2D, ktt::DimensionVector()); - m_triangularToSymmetricDefinition = m_tuner.AddKernelDefinitionFromFile("triangular_to_symmetric", m_kernelFile, ndRangeDim2D, workGroupDim2D); + m_meanDefinition = m_tuner->AddKernelDefinitionFromFile("mean_kernel", m_kernelFile, ndRangeDim1D, workGroupDim1D); + m_reduceDefinition = m_tuner->AddKernelDefinitionFromFile("reduce_kernel", m_kernelFile, ndRangeDim2D, workGroupDim2D); + m_covarDefinition = m_tuner->AddKernelDefinitionFromFile("covar_kernel", m_kernelFile, ndRangeDim1D, workGroupDim1D); + m_gemmDefinition = m_tuner->AddKernelDefinitionFromFile("gemm_fast", m_gemmFile, ndRangeDim2D, ktt::DimensionVector()); + m_triangularToSymmetricDefinition = m_tuner->AddKernelDefinitionFromFile("triangular_to_symmetric", m_kernelFile, ndRangeDim2D, workGroupDim2D); - m_kernel = m_tuner.CreateCompositeKernel("Covariance", + m_kernel = m_tuner->CreateCompositeKernel("Covariance", {m_refMeanDefinition, m_refReduceDefinition, m_refCovarDefinition, m_meanDefinition, m_reduceDefinition, m_covarDefinition, m_gemmDefinition, m_triangularToSymmetricDefinition}, [this](ktt::ComputeInterface& interface) @@ -118,14 +118,14 @@ class Covariance : public ExampleReferenceComputation { } }); - m_tuner.SetArguments(m_refMeanDefinition, {mMeanId, mDataId, mFloatNId, mMId, mNId}); - m_tuner.SetArguments(m_refReduceDefinition, {mMeanId, mDataId, mMId, mNId}); - m_tuner.SetArguments(m_refCovarDefinition, {mSymmatId, mDataId, mMId, mNId}); - m_tuner.SetArguments(m_meanDefinition, {mMeanId, mDataId, mFloatNId, mMId, mNId}); - m_tuner.SetArguments(m_reduceDefinition, {mMeanId, mDataId, mMId, mNId}); - m_tuner.SetArguments(m_covarDefinition, {mSymmatId, mDataId, mMId, mNId}); - m_tuner.SetArguments(m_gemmDefinition, {mMId, mNId, mDataId, mSymmatId}); - m_tuner.SetArguments(m_triangularToSymmetricDefinition, {mSymmatId, mMId}); + m_tuner->SetArguments(m_refMeanDefinition, {mMeanId, mDataId, mFloatNId, mMId, mNId}); + m_tuner->SetArguments(m_refReduceDefinition, {mMeanId, mDataId, mMId, mNId}); + m_tuner->SetArguments(m_refCovarDefinition, {mSymmatId, mDataId, mMId, mNId}); + m_tuner->SetArguments(m_meanDefinition, {mMeanId, mDataId, mFloatNId, mMId, mNId}); + m_tuner->SetArguments(m_reduceDefinition, {mMeanId, mDataId, mMId, mNId}); + m_tuner->SetArguments(m_covarDefinition, {mSymmatId, mDataId, mMId, mNId}); + m_tuner->SetArguments(m_gemmDefinition, {mMId, mNId, mDataId, mSymmatId}); + m_tuner->SetArguments(m_triangularToSymmetricDefinition, {mSymmatId, mMId}); } void InitTuningSpace() override @@ -134,36 +134,36 @@ class Covariance : public ExampleReferenceComputation { // Some parameters are commented out to cut down the tuned space - it is now the same as the // simpler and commonly tuned space in CLBlast (plus our new parameters). // KERNEL: 1 - reference kernels, 0 - edited reference kernels, 2 - use GEMM as third kernel - m_tuner.AddParameter(m_kernel, "KERNEL", vector{1, 0, 2}); - m_tuner.AddParameter(m_kernel, "MWG", vector{16, 32, 64, 128}); - m_tuner.AddParameter(m_kernel, "NWG", vector{16, 32, 64, 128}); - m_tuner.AddParameter(m_kernel, "KWG", vector{16, 32}); - m_tuner.AddParameter(m_kernel, "MDIMC", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "NDIMC", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "MDIMA", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "NDIMB", vector{8, 16, 32}); - m_tuner.AddParameter(m_kernel, "KWI", vector{2, 8}); - m_tuner.AddParameter(m_kernel, "VWM", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "VWN", vector{1, 2, 4, 8}); - m_tuner.AddParameter(m_kernel, "STRM", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "STRN", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "SA", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "SB", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "SYMMETRIC", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "KERNEL", vector{1, 0, 2}); + m_tuner->AddParameter(m_kernel, "MWG", vector{16, 32, 64, 128}); + m_tuner->AddParameter(m_kernel, "NWG", vector{16, 32, 64, 128}); + m_tuner->AddParameter(m_kernel, "KWG", vector{16, 32}); + m_tuner->AddParameter(m_kernel, "MDIMC", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "NDIMC", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "MDIMA", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "NDIMB", vector{8, 16, 32}); + m_tuner->AddParameter(m_kernel, "KWI", vector{2, 8}); + m_tuner->AddParameter(m_kernel, "VWM", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "VWN", vector{1, 2, 4, 8}); + m_tuner->AddParameter(m_kernel, "STRM", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "STRN", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "SA", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "SB", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "SYMMETRIC", vector{0, 1}); // If SYMMETRIC == 1: // SYM_STORE == 1: if VWM == 1, store the symmetric value right in the GEMM kernel // SYM_STORE == 0: use fourth kernel to make the triangular matrix symmetric - m_tuner.AddParameter(m_kernel, "SYM_STORE", vector{0, 1}); - m_tuner.AddParameter(m_kernel, "PRECISION", vector{32}); + m_tuner->AddParameter(m_kernel, "SYM_STORE", vector{0, 1}); + m_tuner->AddParameter(m_kernel, "PRECISION", vector{32}); auto globalModifier = [](const uint64_t size, const vector& v) {return (size * v[0] / v[1]);}; - m_tuner.AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"MDIMC", "MWG"}, globalModifier); - m_tuner.AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"NDIMC", "NWG"}, globalModifier); + m_tuner->AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::X, {"MDIMC", "MWG"}, globalModifier); + m_tuner->AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Global, ktt::ModifierDimension::Y, {"NDIMC", "NWG"}, globalModifier); auto localModifier = [](const uint64_t /*size*/, const vector& v) {return v[0];}; - m_tuner.AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"MDIMC"}, localModifier); - m_tuner.AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"NDIMC"}, localModifier); + m_tuner->AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::X, {"MDIMC"}, localModifier); + m_tuner->AddThreadModifier(m_kernel, {m_gemmDefinition}, ktt::ModifierType::Local, ktt::ModifierDimension::Y, {"NDIMC"}, localModifier); auto multipleOfX = [](const std::vector& v) { return IsMultiple(v[0], v[1]); }; auto multipleOfXMulY = [](const std::vector& v) { return IsMultiple(v[0], v[1] * v[2]); }; @@ -173,19 +173,19 @@ class Covariance : public ExampleReferenceComputation { }; // Sets constraints: Requirement for unrolling the KWG loop - m_tuner.AddConstraint(m_kernel, {"KWG", "KWI"}, multipleOfX); + m_tuner->AddConstraint(m_kernel, {"KWG", "KWI"}, multipleOfX); // Sets constraints: Required for integer MWI and NWI - m_tuner.AddConstraint(m_kernel, {"MWG", "MDIMC", "VWM"}, multipleOfXMulY); - m_tuner.AddConstraint(m_kernel, {"NWG", "NDIMC", "VWN"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"MWG", "MDIMC", "VWM"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"NWG", "NDIMC", "VWN"}, multipleOfXMulY); // Sets constraints: Required for integer MWIA and NWIB - m_tuner.AddConstraint(m_kernel, {"MWG", "MDIMA", "VWM"}, multipleOfXMulY); - m_tuner.AddConstraint(m_kernel, {"NWG", "NDIMB", "VWN"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"MWG", "MDIMA", "VWM"}, multipleOfXMulY); + m_tuner->AddConstraint(m_kernel, {"NWG", "NDIMB", "VWN"}, multipleOfXMulY); // Sets constraints: KWG has to be a multiple of KDIMA = ((MDIMC*NDIMC)/(MDIMA)) and KDIMB = (...) - m_tuner.AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "MDIMA"}, multipleOfXMulYDivZ); - m_tuner.AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "NDIMB"}, multipleOfXMulYDivZ); + m_tuner->AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "MDIMA"}, multipleOfXMulYDivZ); + m_tuner->AddConstraint(m_kernel, {"KWG", "MDIMC", "NDIMC", "NDIMB"}, multipleOfXMulYDivZ); // Don't use parameters for polybench reference kernels, auto reference = [](const vector& v) { @@ -197,12 +197,12 @@ class Covariance : public ExampleReferenceComputation { && v[9] == 1 && v[10] == 1 && v[11] == 0 && v[12] == 0 && v[13] == 1 && v[14] == 1 && v[15] == 0; }; - m_tuner.AddConstraint(m_kernel, {"KERNEL", "MWG", "NWG", "KWG", "MDIMC", "NDIMC", "MDIMA", "NDIMB", "KWI", "VWM", "VWN", "STRM", + m_tuner->AddConstraint(m_kernel, {"KERNEL", "MWG", "NWG", "KWG", "MDIMC", "NDIMC", "MDIMA", "NDIMB", "KWI", "VWM", "VWN", "STRM", "STRN", "SA", "SB", "SYMMETRIC"}, reference); // New NVidia GPUs have max. workgroup size auto maxWgSize = [this](const vector& v) {return v[0] * v[1] <= MAX_WORK_GROUP_SIZE;}; - m_tuner.AddConstraint(m_kernel, {"MDIMC", "NDIMC"}, maxWgSize); + m_tuner->AddConstraint(m_kernel, {"MDIMC", "NDIMC"}, maxWgSize); // Symmetric store can't be used for vectors auto symmetric = [](const vector& v) @@ -218,14 +218,14 @@ class Covariance : public ExampleReferenceComputation { return v[1] == 0; }; - m_tuner.AddConstraint(m_kernel, {"SYMMETRIC", "SYM_STORE", "VWM"}, symmetric); + m_tuner->AddConstraint(m_kernel, {"SYMMETRIC", "SYM_STORE", "VWM"}, symmetric); } void InitReference() override { - m_tuner.SetValidationMethod(ktt::ValidationMethod::SideBySideComparison, 128.0); + m_tuner->SetValidationMethod(ktt::ValidationMethod::SideBySideComparison, 128.0); - m_tuner.SetReferenceComputation(mSymmatId, [this](void* buffer) + m_tuner->SetReferenceComputation(mSymmatId, [this](void* buffer) { const float floatN = static_cast(m_n); vector mean(m_m, 0.0f); @@ -263,7 +263,7 @@ class Covariance : public ExampleReferenceComputation { } }); - m_tuner.SetReferenceComputation(mMeanId, [this](void* buffer) + m_tuner->SetReferenceComputation(mMeanId, [this](void* buffer) { const float floatN = static_cast(m_n); float* mean = static_cast(buffer); diff --git a/Examples/Dummy/Dummy.cpp b/Examples/Dummy/Dummy.cpp index 4335d2a1..4c0e9c50 100644 --- a/Examples/Dummy/Dummy.cpp +++ b/Examples/Dummy/Dummy.cpp @@ -17,19 +17,19 @@ using namespace std; class Dummy : public ExampleBase { protected: - Dummy(shared_ptr config, int defaultProblemSize, + Dummy(int argc, char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName) : - ExampleBase(config, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) + ExampleBase(argc, argv, defaultProblemSize, exampleFolderPath, defaultKernelFileBaseName) { m_gridSize = 256; m_atoms = 1024; - m_tuner.SetTimeUnit(ktt::TimeUnit::Microseconds); + m_tuner->SetTimeUnit(ktt::TimeUnit::Microseconds); UseFastMath(); // Set precise measurement by default. - if (m_config->preciseParams == nullopt) { - m_config->preciseParams = ktt::PreciseMeasurementParameters(2000, 20000, 0.005, ktt::DurationCalculationMethod::Minimum); + if (m_preciseParams == nullopt) { + m_preciseParams = ktt::PreciseMeasurementParameters(2000, 20000, 0.005, ktt::DurationCalculationMethod::Minimum); } } @@ -89,21 +89,21 @@ class Dummy : public ExampleBase { const ktt::DimensionVector ndRangeDimensions(m_gridSize / 32, m_gridSize / 4, m_gridSize); const ktt::DimensionVector workGroupDimensions(32, 4); - m_atomInfoId = m_tuner.AddArgumentVector(m_atomInfo, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoXId = m_tuner.AddArgumentVector(m_atomInfoX, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoYId = m_tuner.AddArgumentVector(m_atomInfoY, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoZId = m_tuner.AddArgumentVector(m_atomInfoZ, ktt::ArgumentAccessType::ReadOnly); - m_atomInfoWId = m_tuner.AddArgumentVector(m_atomInfoW, ktt::ArgumentAccessType::ReadOnly); - m_atomsId = m_tuner.AddArgumentScalar(m_atoms); - m_gridSpacingId = m_tuner.AddArgumentScalar(m_gridSpacing); - m_gridDimId = m_tuner.AddArgumentScalar(m_gridSize); - m_energyGridId = m_tuner.AddArgumentVector(m_energyGrid, ktt::ArgumentAccessType::WriteOnly); + m_atomInfoId = m_tuner->AddArgumentVector(m_atomInfo, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoXId = m_tuner->AddArgumentVector(m_atomInfoX, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoYId = m_tuner->AddArgumentVector(m_atomInfoY, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoZId = m_tuner->AddArgumentVector(m_atomInfoZ, ktt::ArgumentAccessType::ReadOnly); + m_atomInfoWId = m_tuner->AddArgumentVector(m_atomInfoW, ktt::ArgumentAccessType::ReadOnly); + m_atomsId = m_tuner->AddArgumentScalar(m_atoms); + m_gridSpacingId = m_tuner->AddArgumentScalar(m_gridSpacing); + m_gridDimId = m_tuner->AddArgumentScalar(m_gridSize); + m_energyGridId = m_tuner->AddArgumentVector(m_energyGrid, ktt::ArgumentAccessType::WriteOnly); InitKernelDefault("directCoulombSum", "CoulombSum", ndRangeDimensions, {m_atomInfoId, m_atomInfoXId, m_atomInfoYId, m_atomInfoZId, m_atomInfoWId, m_atomsId, m_gridSpacingId, m_gridDimId, m_energyGridId}); - m_tuner.SetLauncher(m_kernel, [this](ktt::ComputeInterface& interface) + m_tuner->SetLauncher(m_kernel, [this](ktt::ComputeInterface& interface) { uint64_t sleep = sleepDuration; if (randomizeSleep) @@ -115,10 +115,10 @@ class Dummy : public ExampleBase { void InitTuningSpace() override { - m_tuner.AddParameter(m_kernel, "DUMMY_1", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); - m_tuner.AddParameter(m_kernel, "DUMMY_2", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); - m_tuner.AddParameter(m_kernel, "DUMMY_3", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); - m_tuner.AddParameter(m_kernel, "DUMMY_4", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); + m_tuner->AddParameter(m_kernel, "DUMMY_1", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); + m_tuner->AddParameter(m_kernel, "DUMMY_2", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); + m_tuner->AddParameter(m_kernel, "DUMMY_3", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); + m_tuner->AddParameter(m_kernel, "DUMMY_4", vector{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}); } }; diff --git a/Examples/ExampleBase.cpp b/Examples/ExampleBase.cpp index 0dd5308c..ef41da7a 100644 --- a/Examples/ExampleBase.cpp +++ b/Examples/ExampleBase.cpp @@ -1,10 +1,11 @@ #include "ExampleBase.h" #include "Api/Output/KernelResult.h" #include "ComputeEngine/ComputeApi.h" -#include "ExampleConfigurator.h" +#include "CliComponent.h" #include "Utility/Logger/Logger.h" #include "Utility/Logger/LoggingLevel.h" #include +#include #include #include #include @@ -31,23 +32,6 @@ string ExampleBase::GetKernelFilePath(string exampleFolderPath, string baseName, return kernelPrefix + exampleFolderPath + "/" + baseName + (suffix == nullopt ?defaultKernelFileSuffix : suffix.value()); } -void ExampleBase::PrintRunStats(const string& phaseName, const RunStats& stats, double throughput) -{ - cout << "\n--- " << phaseName << " complete ---" << endl; - cout << "Total runs: " << stats.totalRuns << endl; - cout << "Successful runs: " << stats.successfulRuns << "/" << stats.totalRuns << endl; - if (!stats.bestConfig.empty()) { - cout << "Best configuration: " << stats.bestConfig << endl; - cout << "Best duration: " << stats.bestDuration << " ns" << endl; - } - if (throughput != -1) - { - cout << "Throughput: " << fixed << setprecision(2) << throughput << " runs/s" << endl; - cout.unsetf(ios_base::floatfield); - cout.precision(6); - } -} - void ExampleBase::PrintProgress(const string& phaseName, int currentRun, double elapsedSeconds, double timeBudget, double bestDuration, double throughput) { @@ -59,20 +43,7 @@ void ExampleBase::PrintProgress(const string& phaseName, int currentRun, double cout.precision(6); } -void ExampleBase::RunStats::Update(ktt::KernelResult result) { - totalRuns++; - - if (result.GetStatus() == ktt::ResultStatus::Ok) { - successfulRuns++; - double duration = result.GetTotalDuration(); - if (duration < bestDuration) { - bestDuration = duration; - bestConfig = result.GetConfiguration().GetString(); - } - } -} - -ExampleBase::RunStats ExampleBase::RunTuningPhase( +RunStats ExampleBase::RunTuningPhase( const std::chrono::steady_clock::time_point& startTime, double timeBudgetSeconds, int printInterval) { RunStats stats; @@ -85,7 +56,7 @@ ExampleBase::RunStats ExampleBase::RunTuningPhase( break; } - const auto result = m_tuner.TuneIteration(m_kernel, {}, false, m_config->preciseParams); + const auto result = m_tuner->TuneIteration(m_kernel, {}, false, m_preciseParams); stats.Update(result); if (stats.totalRuns % printInterval == 0 || stats.totalRuns == 1) { @@ -98,7 +69,7 @@ ExampleBase::RunStats ExampleBase::RunTuningPhase( return stats; } -ExampleBase::RunStats ExampleBase::RunExecutionPhase( +RunStats ExampleBase::RunExecutionPhase( const std::chrono::steady_clock::time_point& startTime, double timeBudgetSeconds, const ktt::KernelConfiguration& bestConfig, int printInterval) { @@ -112,7 +83,7 @@ ExampleBase::RunStats ExampleBase::RunExecutionPhase( break; } - const auto result = m_tuner.Run(m_kernel, bestConfig, {}); + const auto result = m_tuner->Run(m_kernel, bestConfig, {}); stats.Update(result); if (stats.totalRuns % printInterval == 0) { @@ -129,7 +100,7 @@ void ExampleBase::RunDynamic() { ktt::Logger::GetLogger().SetLoggingLevel(ktt::LoggingLevel::Warning); const auto startTime = std::chrono::steady_clock::now(); - const double timeBudgetSeconds = m_config->dynamicTuningTime > 0 ? m_config->dynamicTuningTime : 60.0; + const double timeBudgetSeconds = m_dynamicTuningTime > 0 ? m_dynamicTuningTime : 60.0; constexpr int printInterval = 50; RunStats tuningStats = RunTuningPhase(startTime, timeBudgetSeconds, printInterval); @@ -138,9 +109,11 @@ void ExampleBase::RunDynamic() double tuningElapsed = std::chrono::duration(tuningPhaseEnd - startTime).count(); double tuningThroughput = tuningElapsed > 0 ? tuningStats.totalRuns / tuningElapsed : 0; - PrintRunStats("Tuning phase", tuningStats, tuningThroughput); + tuningStats.Print("Tuning phase", tuningThroughput); + + m_compilerTuning->Run(); - const auto bestConfigData = m_tuner.GetBestConfiguration(m_kernel); + const auto bestConfigData = m_tuner->GetBestConfiguration(m_kernel); cout << "\n--- Running with best configuration ---" << endl; const auto runStartTime = std::chrono::steady_clock::now(); @@ -150,16 +123,16 @@ void ExampleBase::RunDynamic() double runElapsed = std::chrono::duration(totalEndTime - runStartTime).count(); double runThroughput = runElapsed > 0 ? runStats.totalRuns / runElapsed : 0; - PrintRunStats("Final statistics", runStats, runThroughput); + runStats.Print("Final statistics", runThroughput); } void ExampleBase::RunOffline() { const auto startTime = std::chrono::steady_clock::now(); - const auto results = m_tuner.Tune(m_kernel, std::move(m_config->stopCondition), m_config->preciseParams); - m_tuner.SaveResults(results, "Output", ktt::OutputFormat::XML); - m_tuner.SaveResults(results, "Output", ktt::OutputFormat::JSON); + const auto results = m_tuner->Tune(m_kernel, std::move(m_stopCondition), m_preciseParams); + m_tuner->SaveResults(results, "Output", ktt::OutputFormat::XML); + m_tuner->SaveResults(results, "Output", ktt::OutputFormat::JSON); const auto endTime = std::chrono::steady_clock::now(); double elapsed = std::chrono::duration(endTime - startTime).count(); @@ -179,17 +152,20 @@ void ExampleBase::RunOffline() } stats.totalRuns = static_cast(results.size()); - PrintRunStats("Offline tuning", stats, throughput); + stats.Print("Offline tuning", throughput); + + m_compilerTuning->Run(); } void ExampleBase::Run() { - if (m_config->useDynamicTuning) RunDynamic(); + if (m_useDynamicTuning) RunDynamic(); else RunOffline(); } ExampleBase::ExampleBase( - shared_ptr config, + int argc, + char **argv, int defaultProblemSize, string exampleFolderPath, string defaultKernelFileBaseName @@ -201,42 +177,166 @@ ExampleBase::ExampleBase( #elif KTT_CPP_EXAMPLE m_computeApi(ktt::ComputeApi::Cpp), #endif - m_config(config), - m_tuner(config->platform, config->device, m_computeApi) + m_argc(argc), + m_argv(argv), + m_compilerTuning(make_unique(m_tuner, m_kernel)), + m_preheatingSeconds(0) { - m_problemSize = config->problemSize >= 0 ? config->problemSize : defaultProblemSize; - m_kernelFile = config->kernelFile.empty() - ? GetKernelFilePath(exampleFolderPath, defaultKernelFileBaseName) - : config->kernelFile; - - - if (config->useProfiling) - { - printf("Executing with profiling switched ON.\n"); - m_tuner.SetProfiling(true); - } - - m_tuner.SetGlobalSizeType(ktt::GlobalSizeType::CUDA); - m_tuner.SetTimeUnit(ktt::TimeUnit::Microseconds); + m_problemSize = defaultProblemSize; + m_kernelFile = GetKernelFilePath(exampleFolderPath, defaultKernelFileBaseName); } void ExampleBase::PostInitialize() { + InitCLI(); + ProcessCLI(); + InitTuner(); InitData(); InitKernel(); InitTuningSpace(); + Preheat(); InitSearcher(); } +void ExampleBase::InitCLI() { + m_cli.AddOption({[this](const vector &) { + m_rapidTest = true; + }, "--rapidTest", "Run in rapid test mode"}); + + m_cli.AddOption({[this](const vector &) { + m_useProfiling = true; + }, "--profile", "Enable profiling"}); + + m_cli.AddOption({[this](const vector &args) { + m_platform = stoul(args[0]); + }, "--platform", "Platform index (expects int)", "", 1}); + + m_cli.AddOption({[this](const vector &args) { + m_device = stoul(args[0]); + }, "--device", "Device index (expects int)", "", 1}); + + m_cli.AddOption({[this](const vector &args) { + m_problemSize = stoi(args[0]); + }, "--problemSize", "Problem size in MiB (expects int)", "", 1}); + + m_cli.AddOption({[this](const vector &args) { + m_kernelFile = args[0]; + }, "--kernelPath", "Kernel file path (expects string)", "", 1}); + + m_cli.AddOption({[this](const vector &args) { + if (args[0] == "ds") { + m_searcher = make_unique(); + } else if (args[0] == "random") { + m_searcher = make_unique(); + } else if (args[0] == "mcmc") { + m_searcher = make_unique(); + } else { + cerr << "--searcher expects one of (ds, random, mcmc)\n"; + exit(1); + } + }, "--searcher", "Searcher type (ds, random, mcmc)", "", 1}); + + m_cli.AddOption({[this](const vector &args) { + m_profileSearchModelPath = args[0]; + m_useProfiling = true; + }, "--profileSearcher", + "Enable profile searcher and set path to model (expects string) (functions only on CUDA devices)", + "", 1}); + + m_cli.AddOption({[this](const vector &args) { + if (args[0] == "confs") { + m_stopCondition = make_unique(stoul(args[1])); + } else if (args[0] == "fails") { + m_stopCondition = make_unique(stoul(args[1])); + } else if (args[0] == "time") { + m_stopCondition = make_unique(stod(args[1])); + } else if (args[0] == "best") { + m_stopCondition = make_unique(stod(args[1])); + } else { + cerr << "--stopCondition expects one of (confs, fails, time, best)\n"; + exit(1); + } + }, "--stopCondition", + "Set a stop condition. can be confs, fails, time, best. " + " is respectively configuration count (ulong), failed kernel run count (ulong), " + "total tuning duration in seconds (double), best configuration duration in milliseconds (double).", + " ", 2}); + + m_cli.AddOption({[this](const vector &args) { + m_preciseParams = ktt::PreciseMeasurementParameters(stoul(args[0]), + stoul(args[1]), stod(args[2])); + }, "--preciseParams", "Set PreciseMeasurementParameters, calculationDurationMethod is the default Minimum, refer to KTT documentation for details.", + " ", 3}); + m_cli.AddOption({[this](const vector &args) { + if (m_preciseParams == std::nullopt) { + cerr << "--preciseParams must be used before this option.\n"; + exit(1); + } + ktt::DurationCalculationMethod calcMethod = ktt::DurationCalculationMethod::Minimum; + if (args[0] == "min") {} + else if (args[0] == "median") { + calcMethod = ktt::DurationCalculationMethod::Median; + } else if (args[0] == "avg") { + calcMethod = ktt::DurationCalculationMethod::Average; + } else { + cerr << "--preciseParamsCalcMethod expects one of (min, median, avg)\n"; + exit(1); + } + m_preciseParams->durationCalculationMethod = calcMethod; + }, "--preciseParamsCalcMethod", "Optionally set PreciseMeasurementParameters::durationCalculationMethod AFTER USING --preciseParams, expects one of " + "(min, median, avg), refer to KTT documentation for details.", + "", 1}); + + m_cli.AddOption({[this](const vector &args) { + m_preheatingSeconds = stof(args[0]); + }, "--preheat", "Set time in seconds to spend preheating GPU before tuning starts. 0 (disabled) is default. Expects float.", "