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

Minor fix for recent PR (#260)

Browse Source 
* fix example

* update IsSupportedArgument

* fix

* disable fp64 conv example as test

[ROCm/composable_kernel commit: 85fc91c321]
This commit is contained in:
Chao Liu
2022-05-30 19:57:49 -05:00
committed by GitHub
parent 7841891136
commit d56cc51f7d
12 changed files with 62 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
example/01_gemm/CMakeLists.txt
View File

@@ -4,4 +4,5 @@ add_example_executable(example_gemm_dl_int8 gemm_dl_int8.cpp)
add_example_executable(example_gemm_xdl_fp16 gemm_xdl_fp16.cpp)
add_example_executable(example_gemm_xdl_bf16 gemm_xdl_bf16.cpp)
add_example_executable(example_gemm_xdl_int8 gemm_xdl_int8.cpp)
add_example_executable(example_gemm_xdl_fp64 gemm_xdl_fp64.cpp)
# FIXME: re-enable this exampe as test when SWDEV-335738 is fixed
add_example_executable_no_testing(example_gemm_xdl_fp64 gemm_xdl_fp64.cpp)

4
example/01_gemm/gemm_dl_fp16.cpp
View File

@@ -170,9 +170,7 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
std::cout << "wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem"
<< std::endl;
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}

4
example/01_gemm/gemm_dl_fp32.cpp
View File

@@ -169,9 +169,7 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
std::cout << "wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem"
<< std::endl;
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}

4
example/01_gemm/gemm_dl_int8.cpp
View File

@@ -167,9 +167,7 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
std::cout << "wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem"
<< std::endl;
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}

6
example/01_gemm/gemm_xdl_bf16.cpp
View File

@@ -193,9 +193,9 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
throw std::runtime_error(
"wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem");
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}
float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel});

6
example/01_gemm/gemm_xdl_fp16.cpp
View File

@@ -166,9 +166,9 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
throw std::runtime_error(
"wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem");
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}
float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel});

10
example/01_gemm/gemm_xdl_fp64.cpp
View File

@@ -21,8 +21,6 @@ template <ck::index_t... Is>
using S = ck::Sequence<Is...>;
using F64 = double;
using F32 = float;
using F16 = ck::half_t;
using ADataType = double;
using BDataType = double;
@@ -195,9 +193,9 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
throw std::runtime_error(
"wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem");
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}
float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel});
@@ -233,7 +231,7 @@ int main(int argc, char* argv[])
show_2d_matrix(std::cout << "c_host :", c_m_n_host_result) << std::endl;
}
#endif
ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData);
return ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData) ? 0 : 1;
}
return 0;

6
example/01_gemm/gemm_xdl_int8.cpp
View File

@@ -194,9 +194,9 @@ int main(int argc, char* argv[])
if(!gemm.IsSupportedArgument(argument))
{
throw std::runtime_error(
"wrong! device_gemm with the specified compilation parameters does "
"not support this GEMM problem");
std::cout << gemm.GetTypeString() << " does not support this problem" << std::endl;
return 0;
}
float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel});

3
example/09_convnd_fwd/CMakeLists.txt
View File

@@ -1,7 +1,8 @@
add_example_executable(example_convnd_fwd_xdl_fp32 convnd_fwd_xdl_fp32.cpp)
add_example_executable(example_convnd_fwd_xdl_int8 convnd_fwd_xdl_int8.cpp)
add_example_executable(example_convnd_fwd_xdl_fp16 convnd_fwd_xdl_fp16.cpp)
add_example_executable(example_convnd_fwd_xdl_fp64 convnd_fwd_xdl_fp64.cpp)
# FIXME: re-enable this exampe as test when SWDEV-335738 is fixed
add_example_executable_no_testing(example_convnd_fwd_xdl_fp64 convnd_fwd_xdl_fp64.cpp)
target_link_libraries(example_convnd_fwd_xdl_fp64 PRIVATE conv_util)
target_link_libraries(example_convnd_fwd_xdl_fp32 PRIVATE conv_util)
target_link_libraries(example_convnd_fwd_xdl_int8 PRIVATE conv_util)

26
include/ck/tensor_operation/gpu/device/device_convnd_fwd_xdl_nhwc_kyxc_nhwk.hpp
View File

@@ -1,5 +1,4 @@
#ifndef DEVICE_CONVND_FWD_XDL_NHWC_KYXC_NHWK_HPP
#define DEVICE_CONVND_FWD_XDL_NHWC_KYXC_NHWK_HPP
#pragma once
#include <functional>
#include <iostream>
@@ -8,6 +7,7 @@
#include <sstream>
#include "device.hpp"
#include "device_prop.hpp"
#include "device_base.hpp"
#include "device_conv_fwd.hpp"
#include "convolution_forward_specialization.hpp"
@@ -858,6 +858,27 @@ struct DeviceConvNDFwdXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K
static bool IsSupportedArgument(const Argument& arg)
{
if(ck::get_device_name() == "gfx908")
{
if constexpr(!(is_same_v<AccDataType, float> || is_same_v<AccDataType, float> ||
is_same_v<AccDataType, int32_t>))
{
return false;
}
}
else if(ck::get_device_name() == "gfx90a")
{
if constexpr(!(is_same_v<AccDataType, float> || is_same_v<AccDataType, float> ||
is_same_v<AccDataType, int32_t> || is_same_v<AccDataType, double>))
{
return false;
}
}
else
{
return false;
}
// Input tensors can't be bigger than 2GB each.
constexpr ck::long_index_t GB2 = (ck::long_index_t{1} << 31);
@@ -1021,4 +1042,3 @@ struct DeviceConvNDFwdXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K
} // namespace device
} // namespace tensor_operation
} // namespace ck
#endif

2
include/ck/tensor_operation/gpu/device/device_gemm_dl.hpp
View File

@@ -4,6 +4,7 @@
#include <sstream>
#include "device.hpp"
#include "device_prop.hpp"
#include "device_base.hpp"
#include "device_gemm.hpp"
#include "common_header.hpp"
@@ -13,7 +14,6 @@
#include "gemm_specialization.hpp"
#include "element_wise_operation.hpp"
#include "gridwise_gemm_dl_v1r3.hpp"
#include "device_prop.hpp"
namespace ck {
namespace tensor_operation {

20
include/ck/tensor_operation/gpu/device/device_gemm_xdl.hpp
View File

@@ -3,6 +3,7 @@
#include <iostream>
#include <sstream>
#include "device.hpp"
#include "device_prop.hpp"
#include "device_base.hpp"
#include "device_gemm.hpp"
#include "common_header.hpp"
@@ -11,7 +12,6 @@
#include "tensor_descriptor_helper.hpp"
#include "gridwise_gemm_xdlops_v2r3.hpp"
#include "gemm_specialization.hpp"
#include "device_prop.hpp"
namespace ck {
namespace tensor_operation {
@@ -408,7 +408,23 @@ struct DeviceGemmXdl
static bool IsSupportedArgument(const Argument& arg)
{
if(!(ck::get_device_name() == "gfx908" || ck::get_device_name() == "gfx90a"))
if(ck::get_device_name() == "gfx908")
{
if constexpr(!(is_same_v<AccDataType, float> || is_same_v<AccDataType, float> ||
is_same_v<AccDataType, int32_t>))
{
return false;
}
}
else if(ck::get_device_name() == "gfx90a")
{
if constexpr(!(is_same_v<AccDataType, float> || is_same_v<AccDataType, float> ||
is_same_v<AccDataType, int32_t> || is_same_v<AccDataType, double>))
{
return false;
}
}
else
{
return false;
}