Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel_ROCM
Commits
8e0beb65
Commit
8e0beb65
authored
Jun 06, 2023
by
Alan Turner
Browse files
Add unit tests
parent
fcca3307
Changes
5
Expand all
Show whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
398 additions
and
31 deletions
+398
-31
CMakeLists.txt
CMakeLists.txt
+5
-2
library/src/jit_library/src/device_gemm_multiple_d.cpp
library/src/jit_library/src/device_gemm_multiple_d.cpp
+2
-2
test/CMakeLists.txt
test/CMakeLists.txt
+31
-27
test/jit_library/CMakeLists.txt
test/jit_library/CMakeLists.txt
+3
-0
test/jit_library/jit_library.cpp
test/jit_library/jit_library.cpp
+357
-0
No files found.
CMakeLists.txt
View file @
8e0beb65
...
...
@@ -249,6 +249,8 @@ include_directories(BEFORE
${
HIP_INCLUDE_DIRS
}
)
add_custom_target
(
check COMMAND
${
CMAKE_CTEST_COMMAND
}
--output-on-failure -C
${
CMAKE_CFG_INTDIR
}
)
if
(
NOT CK_BUILD_JIT_LIB
)
SET
(
BUILD_DEV ON CACHE BOOL
"BUILD_DEV"
)
if
(
BUILD_DEV
)
...
...
@@ -257,7 +259,7 @@ if (NOT CK_BUILD_JIT_LIB)
endif
()
message
(
"CMAKE_CXX_FLAGS:
${
CMAKE_CXX_FLAGS
}
"
)
add_custom_target
(
check COMMAND
${
CMAKE_CTEST_COMMAND
}
--output-on-failure -C
${
CMAKE_CFG_INTDIR
}
)
file
(
GLOB_RECURSE INSTANCE_FILES
"
${
PROJECT_SOURCE_DIR
}
/*/device_*_instance.cpp"
)
file
(
GLOB dir_list RELATIVE
${
PROJECT_SOURCE_DIR
}
/library/src/tensor_operation_instance/gpu
${
PROJECT_SOURCE_DIR
}
/library/src/tensor_operation_instance/gpu/*
)
...
...
@@ -286,7 +288,6 @@ if (NOT CK_BUILD_JIT_LIB)
add_subdirectory
(
example
)
add_subdirectory
(
test
)
add_subdirectory
(
profiler
)
else
()
...
...
@@ -297,7 +298,9 @@ else()
endif
()
add_subdirectory
(
library
)
add_subdirectory
(
test
)
#Create an interface target for the include only files and call it "composablekernels"
include
(
CMakePackageConfigHelpers
)
...
...
library/src/jit_library/src/device_gemm_multiple_d.cpp
View file @
8e0beb65
...
...
@@ -49,7 +49,7 @@ std::vector<std::string> Problem::GetInstances(const std::string& arch) const
const
bool
quantize
=
ADataType
==
DataType
::
Int8
and
BDataType
==
DataType
::
Int8
;
if
(
get_xdlop_archs
().
find
(
arch
)
!=
get_xdlop_archs
().
end
())
{
instance
::
gemm_add_add_fastgelu_instances
all_instances
{};
ck
::
tensor_operation
::
device
::
instance
::
gemm_add_add_fastgelu_instances
all_instances
{};
if
(
TransA
and
TransB
)
instances
=
all_instances
.
get_col_col_instances
(
quantize
);
else
if
(
TransA
and
not
TransB
)
...
...
@@ -139,7 +139,7 @@ Solution Problem::MakeSolution(std::size_t idx, const std::string& arch) const
std
::
string
Problem
::
GetIncludeHeader
()
const
{
return
instance
::
gemm_add_add_fastgelu_instances
{}.
get_include_header
();
return
ck
::
tensor_operation
::
device
::
instance
::
gemm_add_add_fastgelu_instances
{}.
get_include_header
();
}
std
::
vector
<
Solution
>
Problem
::
GetSolutions
(
const
std
::
string
&
arch
)
const
...
...
test/CMakeLists.txt
View file @
8e0beb65
...
...
@@ -31,33 +31,37 @@ function(add_gtest_executable TEST_NAME)
rocm_install
(
TARGETS
${
TEST_NAME
}
COMPONENT tests
)
endfunction
(
add_gtest_executable TEST_NAME
)
add_subdirectory
(
magic_number_division
)
add_subdirectory
(
space_filling_curve
)
add_subdirectory
(
conv_util
)
add_subdirectory
(
reference_conv_fwd
)
add_subdirectory
(
gemm
)
add_subdirectory
(
gemm_layernorm
)
add_subdirectory
(
gemm_split_k
)
add_subdirectory
(
gemm_reduce
)
add_subdirectory
(
batched_gemm
)
add_subdirectory
(
batched_gemm_reduce
)
add_subdirectory
(
batched_gemm_gemm
)
add_subdirectory
(
batched_gemm_softmax_gemm
)
add_subdirectory
(
batched_gemm_softmax_gemm_permute
)
add_subdirectory
(
grouped_gemm
)
add_subdirectory
(
reduce
)
add_subdirectory
(
convnd_fwd
)
add_subdirectory
(
convnd_bwd_data
)
add_subdirectory
(
grouped_convnd_fwd
)
add_subdirectory
(
grouped_convnd_bwd_weight
)
add_subdirectory
(
block_to_ctile_map
)
add_subdirectory
(
softmax
)
add_subdirectory
(
normalization
)
add_subdirectory
(
data_type
)
add_subdirectory
(
elementwise_normalization
)
add_subdirectory
(
batchnorm
)
add_subdirectory
(
contraction
)
add_subdirectory
(
pool_fwd
)
if
(
CK_BUILD_JIT_LIB
)
add_subdirectory
(
jit_library
)
else
()
add_subdirectory
(
magic_number_division
)
add_subdirectory
(
space_filling_curve
)
add_subdirectory
(
conv_util
)
add_subdirectory
(
reference_conv_fwd
)
add_subdirectory
(
gemm
)
add_subdirectory
(
gemm_layernorm
)
add_subdirectory
(
gemm_split_k
)
add_subdirectory
(
gemm_reduce
)
add_subdirectory
(
batched_gemm
)
add_subdirectory
(
batched_gemm_reduce
)
add_subdirectory
(
batched_gemm_gemm
)
add_subdirectory
(
batched_gemm_softmax_gemm
)
add_subdirectory
(
batched_gemm_softmax_gemm_permute
)
add_subdirectory
(
grouped_gemm
)
add_subdirectory
(
reduce
)
add_subdirectory
(
convnd_fwd
)
add_subdirectory
(
convnd_bwd_data
)
add_subdirectory
(
grouped_convnd_fwd
)
add_subdirectory
(
grouped_convnd_bwd_weight
)
add_subdirectory
(
block_to_ctile_map
)
add_subdirectory
(
softmax
)
add_subdirectory
(
normalization
)
add_subdirectory
(
data_type
)
add_subdirectory
(
elementwise_normalization
)
add_subdirectory
(
batchnorm
)
add_subdirectory
(
contraction
)
add_subdirectory
(
pool_fwd
)
endif
()
if
(
GPU_TARGETS MATCHES
"gfx1100"
)
add_subdirectory
(
wmma_op
)
endif
()
test/jit_library/CMakeLists.txt
0 → 100644
View file @
8e0beb65
add_test_executable
(
test_jit_library jit_library.cpp
)
add_dependencies
(
test_jit_library jit_library
)
target_link_libraries
(
test_jit_library PRIVATE jit_library
)
test/jit_library/jit_library.cpp
0 → 100644
View file @
8e0beb65
This diff is collapsed.
Click to expand it.
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment