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
f084651e
Commit
f084651e
authored
Apr 26, 2023
by
ltqin
Browse files
fix enum
parent
185af92b
Changes
4
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
13 additions
and
9 deletions
+13
-9
library/include/ck/library/tensor_operation_instance/add_device_operation_instance.hpp
...nsor_operation_instance/add_device_operation_instance.hpp
+5
-3
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/batched_gemm_softmax_gemm_permute.hpp
...oftmax_gemm_permute/batched_gemm_softmax_gemm_permute.hpp
+2
-2
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_multiple_d_softmax_gemm_permute_xdl_cshuffle_bf16_gmk_gnk_gno_gmo_instance.hpp
...mm_permute_xdl_cshuffle_bf16_gmk_gnk_gno_gmo_instance.hpp
+3
-2
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_multiple_d_softmax_gemm_permute_xdl_cshuffle_fp16_gmk_gnk_gno_gmo_instance.hpp
...mm_permute_xdl_cshuffle_fp16_gmk_gnk_gno_gmo_instance.hpp
+3
-2
No files found.
library/include/ck/library/tensor_operation_instance/add_device_operation_instance.hpp
View file @
f084651e
...
@@ -34,14 +34,16 @@ enum struct ArchitectureEnum
...
@@ -34,14 +34,16 @@ enum struct ArchitectureEnum
None
,
None
,
Gfx908
,
Gfx908
,
Gfx90a
,
Gfx90a
,
Gfx940
,
Gfx1030
,
Gfx1030
,
All
All
};
};
enum
struct
Gemm
FeatureEnum
enum
struct
Arch
FeatureEnum
{
{
None
,
Xdl
,
Xdl
,
Dl
,
Dl
,
Wm
d
Wm
ma
};
};
template
<
ArchitectureEnum
...
Is
>
template
<
ArchitectureEnum
...
Is
>
struct
ArchitectureEnumSequence
struct
ArchitectureEnumSequence
...
@@ -55,7 +57,7 @@ struct ArchitectureEnumSequence
...
@@ -55,7 +57,7 @@ struct ArchitectureEnumSequence
return
mData
[
I
];
return
mData
[
I
];
}
}
};
};
template
<
Gemm
FeatureEnum
Feature
,
typename
DeviceOp
>
template
<
Arch
FeatureEnum
Feature
,
typename
DeviceOp
>
struct
DeviceOperationInstances
;
struct
DeviceOperationInstances
;
template
<
typename
Arch
,
typename
DeviceOp
>
template
<
typename
Arch
,
typename
DeviceOp
>
...
...
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/batched_gemm_softmax_gemm_permute.hpp
View file @
f084651e
...
@@ -74,11 +74,11 @@ struct DeviceOperationInstanceCreator<Arch,
...
@@ -74,11 +74,11 @@ struct DeviceOperationInstanceCreator<Arch,
MaskingSpec
>
;
MaskingSpec
>
;
static
void
add_device_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceOp
>>&
instances
)
static
void
add_device_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceOp
>>&
instances
)
{
{
if
constexpr
(
DeviceOperationInstances
<
Gemm
FeatureEnum
::
Xdl
,
if
constexpr
(
DeviceOperationInstances
<
Arch
FeatureEnum
::
Xdl
,
DeviceOp
>::
template
is_surport
<
Arch
>())
DeviceOp
>::
template
is_surport
<
Arch
>())
add_device_operation_instances
(
add_device_operation_instances
(
instances
,
instances
,
DeviceOperationInstances
<
Gemm
FeatureEnum
::
Xdl
,
DeviceOp
>::
get_device_instances
());
DeviceOperationInstances
<
Arch
FeatureEnum
::
Xdl
,
DeviceOp
>::
get_device_instances
());
}
}
};
};
...
...
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_multiple_d_softmax_gemm_permute_xdl_cshuffle_bf16_gmk_gnk_gno_gmo_instance.hpp
View file @
f084651e
...
@@ -27,7 +27,7 @@ template <index_t NumDimG,
...
@@ -27,7 +27,7 @@ template <index_t NumDimG,
typename
B1ElementwiseOperation
,
typename
B1ElementwiseOperation
,
typename
C1DEElementwiseOperation
,
typename
C1DEElementwiseOperation
,
MaskingSpecialization
MaskingSpec
>
MaskingSpecialization
MaskingSpec
>
struct
DeviceOperationInstances
<
Gemm
FeatureEnum
::
Xdl
,
struct
DeviceOperationInstances
<
Arch
FeatureEnum
::
Xdl
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
NumDimM
,
NumDimM
,
NumDimN
,
NumDimN
,
...
@@ -95,7 +95,8 @@ struct DeviceOperationInstances<GemmFeatureEnum::Xdl,
...
@@ -95,7 +95,8 @@ struct DeviceOperationInstances<GemmFeatureEnum::Xdl,
ck
::
static_for
<
0
,
Archs
::
mSize
,
1
>
{}([
&
](
auto
I
)
{
ck
::
static_for
<
0
,
Archs
::
mSize
,
1
>
{}([
&
](
auto
I
)
{
if
constexpr
(
Archs
::
At
(
I
)
==
ArchitectureEnum
::
All
||
if
constexpr
(
Archs
::
At
(
I
)
==
ArchitectureEnum
::
All
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx908
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx908
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx90a
)
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx90a
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx940
)
is_surport
=
true
;
is_surport
=
true
;
});
});
return
is_surport
;
return
is_surport
;
...
...
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_multiple_d_softmax_gemm_permute_xdl_cshuffle_fp16_gmk_gnk_gno_gmo_instance.hpp
View file @
f084651e
...
@@ -27,7 +27,7 @@ template <index_t NumDimG,
...
@@ -27,7 +27,7 @@ template <index_t NumDimG,
typename
B1ElementwiseOperation
,
typename
B1ElementwiseOperation
,
typename
C1DEElementwiseOperation
,
typename
C1DEElementwiseOperation
,
MaskingSpecialization
MaskingSpec
>
MaskingSpecialization
MaskingSpec
>
struct
DeviceOperationInstances
<
Gemm
FeatureEnum
::
Xdl
,
struct
DeviceOperationInstances
<
Arch
FeatureEnum
::
Xdl
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
NumDimM
,
NumDimM
,
NumDimN
,
NumDimN
,
...
@@ -96,7 +96,8 @@ struct DeviceOperationInstances<GemmFeatureEnum::Xdl,
...
@@ -96,7 +96,8 @@ struct DeviceOperationInstances<GemmFeatureEnum::Xdl,
ck
::
static_for
<
0
,
Archs
::
mSize
,
1
>
{}([
&
](
auto
I
)
{
ck
::
static_for
<
0
,
Archs
::
mSize
,
1
>
{}([
&
](
auto
I
)
{
if
constexpr
(
Archs
::
At
(
I
)
==
ArchitectureEnum
::
All
||
if
constexpr
(
Archs
::
At
(
I
)
==
ArchitectureEnum
::
All
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx908
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx908
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx90a
)
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx90a
||
Archs
::
At
(
I
)
==
ArchitectureEnum
::
Gfx940
)
is_surport
=
true
;
is_surport
=
true
;
});
});
return
is_surport
;
return
is_surport
;
...
...
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