ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-26 08:00:13 +00:00
Code Issues Packages Projects Releases Wiki Activity

[CK][EXAMPLES] (#2826)

Browse Source 
-Added parameter to enable/disable verification and timing of kernel in various examples that missed it.
-Added parameter to change number of groups to execute in grouped_gemm_examples.

Signed-off-by: Michal Kulikowski <Michal.Kulikowski@amd.com>

[ROCm/composable_kernel commit: ffe9775e70]
This commit is contained in:
Michał Kulikowski
2025-09-11 21:33:00 +02:00
committed by GitHub
parent 6e774b512a
commit 064eb037db
33 changed files with 578 additions and 116 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_relu_perchannel_quantization_int8.cpp
View File

@@ -78,8 +78,28 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_bias_perchannel_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
const auto out_element_op = OutElementOp{ActivationOp{}};
run_conv2d_fwd_bias_perchannel_quantization_example(out_element_op);
run_conv2d_fwd_bias_perchannel_quantization_example(
out_element_op, do_verification, time_kernel);
};

23
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_relu_perlayer_quantization_int8.cpp
View File

@@ -76,9 +76,28 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_bias_perlayer_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
float requant_scale = 0.5f;
const auto out_element_op = OutElementOp{requant_scale, ActivationOp{}};
run_conv2d_fwd_bias_perlayer_quantization_example(out_element_op);
run_conv2d_fwd_bias_perlayer_quantization_example(out_element_op, do_verification, time_kernel);
}

24
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_tanh_perchannel_quantization_int8.cpp
View File

@@ -79,9 +79,29 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_bias_perchannel_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
float scale_z_inv = 0.5f;
const auto out_element_op = OutElementOp{scale_z_inv, ActivationOp{}};
run_conv2d_fwd_bias_perchannel_quantization_example(out_element_op);
run_conv2d_fwd_bias_perchannel_quantization_example(
out_element_op, do_verification, time_kernel);
};

23
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_bias_tanh_perlayer_quantization_int8.cpp
View File

@@ -76,10 +76,29 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_bias_perlayer_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
float scale_acc = 0.5f;
float scale_z_inv = 0.5f;
const auto out_element_op = OutElementOp{scale_z_inv, scale_acc, ActivationOp{}};
run_conv2d_fwd_bias_perlayer_quantization_example(out_element_op);
run_conv2d_fwd_bias_perlayer_quantization_example(out_element_op, do_verification, time_kernel);
}

23
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_perchannel_quantization_int8.cpp
View File

@@ -76,8 +76,27 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_perchannel_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
const auto out_element_op = OutElementOp{ActivationOp{}};
run_conv2d_fwd_perchannel_quantization_example(out_element_op);
run_conv2d_fwd_perchannel_quantization_example(out_element_op, do_verification, time_kernel);
}

23
example/40_conv2d_fwd_quantization/conv2d_fwd_dl_perlayer_quantization_int8.cpp
View File

@@ -71,9 +71,28 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_perlayer_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = false;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
float requant_scale = 0.5f;
const auto out_element_op = OutElementOp{requant_scale, ActivationOp{}};
run_conv2d_fwd_perlayer_quantization_example(out_element_op);
run_conv2d_fwd_perlayer_quantization_example(out_element_op, do_verification, time_kernel);
}

24
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_bias_relu_perchannel_quantization_int8.cpp
View File

@@ -82,8 +82,28 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_bias_perchannel_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
const auto out_element_op = OutElementOp{ActivationOp{}};
run_conv2d_fwd_bias_perchannel_quantization_example(out_element_op);
run_conv2d_fwd_bias_perchannel_quantization_example(
out_element_op, do_verification, time_kernel);
};

23
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_bias_relu_perlayer_quantization_int8.cpp
View File

@@ -80,9 +80,28 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_bias_perlayer_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
float requant_scale = 0.5f;
const auto out_element_op = OutElementOp{requant_scale, ActivationOp{}};
run_conv2d_fwd_bias_perlayer_quantization_example(out_element_op);
run_conv2d_fwd_bias_perlayer_quantization_example(out_element_op, do_verification, time_kernel);
}

23
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_perchannel_quantization_int8.cpp
View File

@@ -80,8 +80,27 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_perchannel_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = true;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
const auto out_element_op = OutElementOp{ActivationOp{}};
run_conv2d_fwd_perchannel_quantization_example(out_element_op);
run_conv2d_fwd_perchannel_quantization_example(out_element_op, do_verification, time_kernel);
}

23
example/40_conv2d_fwd_quantization/conv2d_fwd_xdl_perlayer_quantization_int8.cpp
View File

@@ -75,9 +75,28 @@ using DeviceGroupedConvNDFwdInstance =
#include "run_conv2d_fwd_perlayer_quantization_example.inc"
int main()
int main(int argc, char* argv[])
{
bool do_verification = true;
bool time_kernel = false;
if(argc == 1)
{
// use default
}
else if(argc == 3)
{
do_verification = std::stoi(argv[1]);
time_kernel = std::stoi(argv[2]);
}
else
{
printf("arg1: verification (0=no, 1=yes)\n");
printf("arg2: time kernel (0=no, 1=yes)\n");
exit(0);
}
float requant_scale = 0.5f;
const auto out_element_op = OutElementOp{requant_scale, ActivationOp{}};
run_conv2d_fwd_perlayer_quantization_example(out_element_op);
run_conv2d_fwd_perlayer_quantization_example(out_element_op, do_verification, time_kernel);
}

6
example/40_conv2d_fwd_quantization/run_conv2d_fwd_bias_perchannel_quantization_example.inc
View File

@@ -167,10 +167,10 @@ bool run_grouped_conv_fwd(bool do_verification,
return (pass ? 0 : 1);
}
int run_conv2d_fwd_bias_perchannel_quantization_example(const OutElementOp& out_element_op)
int run_conv2d_fwd_bias_perchannel_quantization_example(const OutElementOp& out_element_op,
bool do_verification,
bool time_kernel)
{
bool do_verification = true;
bool time_kernel = true;
const ck::index_t ndim_spatial = 2;
ck::utils::conv::ConvParam conv_param{

6
example/40_conv2d_fwd_quantization/run_conv2d_fwd_bias_perlayer_quantization_example.inc
View File

@@ -155,10 +155,10 @@ bool run_grouped_conv_fwd(bool do_verification,
return (pass ? 0 : 1);
}
int run_conv2d_fwd_bias_perlayer_quantization_example(const OutElementOp& out_element_op)
int run_conv2d_fwd_bias_perlayer_quantization_example(const OutElementOp& out_element_op,
bool do_verification,
bool time_kernel)
{
bool do_verification = true;
bool time_kernel = true;
const ck::index_t ndim_spatial = 2;
ck::utils::conv::ConvParam conv_param{

6
example/40_conv2d_fwd_quantization/run_conv2d_fwd_perchannel_quantization_example.inc
View File

@@ -157,10 +157,10 @@ bool run_grouped_conv_fwd(bool do_verification,
return (pass ? 0 : 1);
}
int run_conv2d_fwd_perchannel_quantization_example(const OutElementOp& out_element_op)
int run_conv2d_fwd_perchannel_quantization_example(const OutElementOp& out_element_op,
bool do_verification,
bool time_kernel)
{
bool do_verification = true;
bool time_kernel = true;
const ck::index_t ndim_spatial = 2;
ck::utils::conv::ConvParam conv_param{

6
example/40_conv2d_fwd_quantization/run_conv2d_fwd_perlayer_quantization_example.inc
View File

@@ -139,10 +139,10 @@ bool run_grouped_conv_fwd(bool do_verification,
return (pass ? 0 : 1);
}
int run_conv2d_fwd_perlayer_quantization_example(const OutElementOp& out_element_op)
int run_conv2d_fwd_perlayer_quantization_example(const OutElementOp& out_element_op,
bool do_verification,
bool time_kernel)
{
bool do_verification = true;
bool time_kernel = false;
const ck::index_t ndim_spatial = 2;
ck::utils::conv::ConvParam conv_param{