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
Commits
520d2122
"src/include/gridwise_convolution_wrapper.hip.hpp" did not exist on "bdbc0eaad175cd4054105cfb3fc812a8526e3b49"
Unverified
Commit
520d2122
authored
Nov 01, 2023
by
Bartłomiej Kocot
Committed by
GitHub
Nov 01, 2023
Browse files
Merge branch 'develop' into dfeng/fwd_f8_fix
parents
f7a43573
306fd506
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
36 additions
and
17 deletions
+36
-17
Jenkinsfile
Jenkinsfile
+30
-14
include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl.hpp
...e/ck/tensor_operation/gpu/device/impl/device_gemm_xdl.hpp
+2
-1
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
...ary/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
+4
-2
No files found.
Jenkinsfile
View file @
520d2122
...
@@ -139,7 +139,7 @@ def buildDocker(install_prefix){
...
@@ -139,7 +139,7 @@ def buildDocker(install_prefix){
else
{
else
{
echo
"Checking for image: ${image_name}"
echo
"Checking for image: ${image_name}"
sh
"docker manifest inspect --insecure ${image_name}"
sh
"docker manifest inspect --insecure ${image_name}"
echo
"Image: ${image_name} found!
!
Skipping building image"
echo
"Image: ${image_name} found! Skipping building image"
}
}
}
}
catch
(
Exception
ex
){
catch
(
Exception
ex
){
...
@@ -213,8 +213,10 @@ def cmake_build(Map conf=[:]){
...
@@ -213,8 +213,10 @@ def cmake_build(Map conf=[:]){
if
(
setup_args
.
contains
(
"gfx94"
)){
if
(
setup_args
.
contains
(
"gfx94"
)){
invocation_tag
=
"gfx94"
invocation_tag
=
"gfx94"
}
}
echo
"invocation tag: ${invocation_tag}"
def
redis_pre_setup_cmd
=
pre_setup_cmd
if
(
check_host
()
&&
params
.
USE_SCCACHE
&&
"${env.CK_SCCACHE}"
!=
"null"
&&
"${invocation_tag}"
!=
""
)
{
if
(
check_host
()
&&
params
.
USE_SCCACHE
&&
"${env.CK_SCCACHE}"
!=
"null"
&&
"${invocation_tag}"
!=
""
)
{
pre_setup_cmd
=
pre_setup_cmd
+
"""
redis_
pre_setup_cmd
=
pre_setup_cmd
+
"""
#!/bin/bash
#!/bin/bash
export ROCM_PATH=/opt/rocm
export ROCM_PATH=/opt/rocm
export SCCACHE_ENABLED=true
export SCCACHE_ENABLED=true
...
@@ -228,18 +230,30 @@ def cmake_build(Map conf=[:]){
...
@@ -228,18 +230,30 @@ def cmake_build(Map conf=[:]){
export SCCACHE_C_CUSTOM_CACHE_BUSTER="${invocation_tag}"
export SCCACHE_C_CUSTOM_CACHE_BUSTER="${invocation_tag}"
echo \$SCCACHE_C_CUSTOM_CACHE_BUSTER
echo \$SCCACHE_C_CUSTOM_CACHE_BUSTER
stunnel ../script/redis-cli.conf
stunnel ../script/redis-cli.conf
(
../script/sccache_wrapper.sh --enforce_redis
set -e
../script/sccache_wrapper.sh --enforce_redis
)
error_code=\$?
if [ \$error_code -ne 0 ]; then
echo "could not connect to the redis server. using sccache locally."
../script/sccache_wrapper.sh
fi
"""
"""
setup_args
=
" -DCMAKE_CXX_COMPILER_LAUNCHER=sccache -DCMAKE_C_COMPILER_LAUNCHER=sccache "
+
setup_args
try
{
def
cmd1
=
conf
.
get
(
"cmd1"
,
"""
${redis_pre_setup_cmd}
"""
)
sh
cmd1
setup_args
=
" -DCMAKE_CXX_COMPILER_LAUNCHER=sccache -DCMAKE_C_COMPILER_LAUNCHER=sccache "
+
setup_args
}
catch
(
Exception
err
){
echo
"could not connect to redis server: ${err.getMessage()}. will not use sccache."
def
cmd2
=
conf
.
get
(
"cmd2"
,
"""
${pre_setup_cmd}
"""
)
sh
cmd2
}
}
else
{
def
cmd3
=
conf
.
get
(
"cmd3"
,
"""
${pre_setup_cmd}
"""
)
sh
cmd3
}
}
def
setup_cmd
=
conf
.
get
(
"setup_cmd"
,
"${cmake_envs} cmake ${setup_args} .. "
)
def
setup_cmd
=
conf
.
get
(
"setup_cmd"
,
"${cmake_envs} cmake ${setup_args} .. "
)
// reduce parallelism when compiling, clang uses too much memory
// reduce parallelism when compiling, clang uses too much memory
def
nt
=
nthreads
()
def
nt
=
nthreads
()
...
@@ -247,14 +261,16 @@ def cmake_build(Map conf=[:]){
...
@@ -247,14 +261,16 @@ def cmake_build(Map conf=[:]){
def
execute_cmd
=
conf
.
get
(
"execute_cmd"
,
""
)
def
execute_cmd
=
conf
.
get
(
"execute_cmd"
,
""
)
def
cmd
=
conf
.
get
(
"cmd"
,
"""
def
cmd
=
conf
.
get
(
"cmd"
,
"""
${pre_setup_cmd}
${setup_cmd}
${setup_cmd}
${build_cmd}
${build_cmd}
${execute_cmd}
${execute_cmd}
"""
)
"""
)
echo
cmd
echo
cmd
sh
cmd
dir
(
"build"
){
sh
cmd
}
// Only archive from master or develop
// Only archive from master or develop
if
(
package_build
==
true
&&
(
env
.
BRANCH_NAME
==
"develop"
||
env
.
BRANCH_NAME
==
"amd-master"
))
{
if
(
package_build
==
true
&&
(
env
.
BRANCH_NAME
==
"develop"
||
env
.
BRANCH_NAME
==
"amd-master"
))
{
...
...
include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl.hpp
View file @
520d2122
...
@@ -184,7 +184,8 @@ struct DeviceGemmXdl : public DeviceGemm<ALayout,
...
@@ -184,7 +184,8 @@ struct DeviceGemmXdl : public DeviceGemm<ALayout,
return
false
;
return
false
;
}
}
}
}
else
if
(
ck
::
get_device_name
()
==
"gfx90a"
||
ck
::
get_device_name
()
==
"gfx940"
)
else
if
(
ck
::
get_device_name
()
==
"gfx90a"
||
ck
::
get_device_name
()
==
"gfx940"
||
ck
::
get_device_name
()
==
"gfx941"
||
ck
::
get_device_name
()
==
"gfx942"
)
{
{
if
constexpr
(
!
(
is_same_v
<
AccDataType
,
float
>
||
is_same_v
<
AccDataType
,
float
>
||
if
constexpr
(
!
(
is_same_v
<
AccDataType
,
float
>
||
is_same_v
<
AccDataType
,
float
>
||
is_same_v
<
AccDataType
,
int32_t
>
||
is_same_v
<
AccDataType
,
double
>
))
is_same_v
<
AccDataType
,
int32_t
>
||
is_same_v
<
AccDataType
,
double
>
))
...
...
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
View file @
520d2122
...
@@ -108,13 +108,15 @@ if (ENABLE_PIPELINE_V2_OPT)
...
@@ -108,13 +108,15 @@ if (ENABLE_PIPELINE_V2_OPT)
CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT=1
CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT=1
)
)
# TODO: The "-vectorize-slp=false" LLVM option is a workaround to prevent inefficient instruction scheduling
# caused by the SLP Vectorizer. Remove this option after fix the SLP Vectorizer issue.
# layout=NT
# layout=NT
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16/km_kn_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16/km_kn_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
COMPILE_OPTIONS
";
;
"
COMPILE_OPTIONS
";
-mllvm;-vectorize-slp=false
"
COMPILE_DEFINITIONS
"
${
WAVES_PER_EU_DEFS
}
;
${
IGLP_OPT_DEFS
}
"
)
COMPILE_DEFINITIONS
"
${
WAVES_PER_EU_DEFS
}
;
${
IGLP_OPT_DEFS
}
"
)
# layout=NN
# layout=NN
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16/km_nk_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16/km_nk_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
COMPILE_OPTIONS
";
;
"
COMPILE_OPTIONS
";
-mllvm;-vectorize-slp=false
"
COMPILE_DEFINITIONS
"
${
WAVES_PER_EU_DEFS
}
;
${
IGLP_OPT_DEFS
}
"
)
COMPILE_DEFINITIONS
"
${
WAVES_PER_EU_DEFS
}
;
${
IGLP_OPT_DEFS
}
"
)
# layout=TT
# layout=TT
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16/mk_kn_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16/mk_kn_mn_default_pipeline_v2_opt_instance.cpp PROPERTIES
...
...
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