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
5f4a0f73
Commit
5f4a0f73
authored
Apr 25, 2023
by
ltqin
Browse files
change arch position
parent
508b643f
Changes
9
Hide whitespace changes
Inline
Side-by-side
Showing
9 changed files
with
36 additions
and
36 deletions
+36
-36
library/include/ck/library/tensor_operation_instance/add_device_operation_instance.hpp
...nsor_operation_instance/add_device_operation_instance.hpp
+2
-2
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
+9
-9
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
-3
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
-3
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_bias_softmax_gemm_permute_xdl_cshuffle_bf16_bf16_bf16_bf16_gmk_gnk_gno_gmo_instance.cpp
...cshuffle_bf16_bf16_bf16_bf16_gmk_gnk_gno_gmo_instance.cpp
+4
-4
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_bias_softmax_gemm_permute_xdl_cshuffle_f16_f16_f16_f16_gmk_gnk_gno_gmo_instance.cpp
...xdl_cshuffle_f16_f16_f16_f16_gmk_gnk_gno_gmo_instance.cpp
+4
-4
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_bf16_bf16_bf16_bf16_gmk_gnk_gno_gmo_instance.cpp
...cshuffle_bf16_bf16_bf16_bf16_gmk_gnk_gno_gmo_instance.cpp
+4
-4
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_f16_f16_f16_f16_gmk_gnk_gno_gmo_instance.cpp
...xdl_cshuffle_f16_f16_f16_f16_gmk_gnk_gno_gmo_instance.cpp
+4
-4
src_example/01_gemm_bias_softmax_gemm_permute/gemm_bias_softmax_gemm_permute.cpp
...s_softmax_gemm_permute/gemm_bias_softmax_gemm_permute.cpp
+3
-3
No files found.
library/include/ck/library/tensor_operation_instance/add_device_operation_instance.hpp
View file @
5f4a0f73
...
...
@@ -55,10 +55,10 @@ struct ArchitectureEnumSequence
return
mData
[
I
];
}
};
template
<
typename
DeviceOp
,
GemmFeatureEnum
Feature
>
template
<
GemmFeatureEnum
Feature
,
typename
DeviceOp
>
struct
DeviceOperationInstances
;
template
<
typename
DeviceOp
,
typename
Arch
>
template
<
typename
Arch
,
typename
DeviceOp
>
struct
DeviceOperationInstanceCreator
;
}
// namespace instance
}
// namespace device
...
...
library/include/ck/library/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/batched_gemm_softmax_gemm_permute.hpp
View file @
5f4a0f73
...
...
@@ -18,7 +18,8 @@ namespace tensor_operation {
namespace
device
{
namespace
instance
{
template
<
index_t
NumDimG
,
template
<
typename
Arch
,
index_t
NumDimG
,
index_t
NumDimM
,
index_t
NumDimN
,
index_t
NumDimK
,
...
...
@@ -34,9 +35,9 @@ template <index_t NumDimG,
typename
C0DEElementwiseOperation
,
typename
B1ElementwiseOperation
,
typename
C1DEElementwiseOperation
,
MaskingSpecialization
MaskingSpec
,
typename
Arch
>
struct
DeviceOperationInstanceCreator
<
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
MaskingSpecialization
MaskingSpec
>
struct
DeviceOperationInstanceCreator
<
Arch
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
NumDimM
,
NumDimN
,
NumDimK
,
...
...
@@ -52,8 +53,7 @@ struct DeviceOperationInstanceCreator<DeviceBatchedGemmSoftmaxGemmPermute<NumDim
C0DEElementwiseOperation
,
B1ElementwiseOperation
,
C1DEElementwiseOperation
,
MaskingSpec
>
,
Arch
>
MaskingSpec
>>
{
using
DeviceOp
=
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
NumDimM
,
...
...
@@ -74,11 +74,11 @@ struct DeviceOperationInstanceCreator<DeviceBatchedGemmSoftmaxGemmPermute<NumDim
MaskingSpec
>
;
static
void
add_device_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceOp
>>&
instances
)
{
if
constexpr
(
DeviceOperationInstances
<
DeviceOp
,
GemmFeatureEnum
::
Xdl
>::
template
is_surport
<
Arch
>())
if
constexpr
(
DeviceOperationInstances
<
GemmFeatureEnum
::
Xdl
,
DeviceOp
>::
template
is_surport
<
Arch
>())
add_device_operation_instances
(
instances
,
DeviceOperationInstances
<
DeviceOp
,
GemmFeatureEnum
::
Xdl
>::
get_device_instances
());
DeviceOperationInstances
<
GemmFeatureEnum
::
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 @
5f4a0f73
...
...
@@ -27,7 +27,8 @@ template <index_t NumDimG,
typename
B1ElementwiseOperation
,
typename
C1DEElementwiseOperation
,
MaskingSpecialization
MaskingSpec
>
struct
DeviceOperationInstances
<
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
struct
DeviceOperationInstances
<
GemmFeatureEnum
::
Xdl
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
NumDimM
,
NumDimN
,
NumDimK
,
...
...
@@ -43,8 +44,7 @@ struct DeviceOperationInstances<DeviceBatchedGemmSoftmaxGemmPermute<NumDimG,
C0DEElementwiseOperation
,
B1ElementwiseOperation
,
C1DEElementwiseOperation
,
MaskingSpec
>
,
GemmFeatureEnum
::
Xdl
>
MaskingSpec
>>
{
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
...
...
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 @
5f4a0f73
...
...
@@ -27,7 +27,8 @@ template <index_t NumDimG,
typename
B1ElementwiseOperation
,
typename
C1DEElementwiseOperation
,
MaskingSpecialization
MaskingSpec
>
struct
DeviceOperationInstances
<
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
struct
DeviceOperationInstances
<
GemmFeatureEnum
::
Xdl
,
DeviceBatchedGemmSoftmaxGemmPermute
<
NumDimG
,
NumDimM
,
NumDimN
,
NumDimK
,
...
...
@@ -43,8 +44,7 @@ struct DeviceOperationInstances<DeviceBatchedGemmSoftmaxGemmPermute<NumDimG,
C0DEElementwiseOperation
,
B1ElementwiseOperation
,
C1DEElementwiseOperation
,
MaskingSpec
>
,
GemmFeatureEnum
::
Xdl
>
MaskingSpec
>>
{
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
...
...
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_bias_softmax_gemm_permute_xdl_cshuffle_bf16_bf16_bf16_bf16_gmk_gnk_gno_gmo_instance.cpp
View file @
5f4a0f73
...
...
@@ -66,8 +66,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
MaskingSpecialization
::
MaskOutUpperTriangle
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
void
add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_instances
(
...
...
@@ -108,8 +108,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskDisabled
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
}
// namespace instance
...
...
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_bias_softmax_gemm_permute_xdl_cshuffle_f16_f16_f16_f16_gmk_gnk_gno_gmo_instance.cpp
View file @
5f4a0f73
...
...
@@ -65,8 +65,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskOutUpperTriangle
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
void
add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_instances
(
...
...
@@ -107,8 +107,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskDisabled
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
}
// namespace instance
...
...
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_bf16_bf16_bf16_bf16_gmk_gnk_gno_gmo_instance.cpp
View file @
5f4a0f73
...
...
@@ -65,8 +65,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskOutUpperTriangle
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
void
add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_instances
(
...
...
@@ -107,8 +107,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskDisabled
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
}
// namespace instance
...
...
library/src/tensor_operation_instance/gpu/batched_gemm_softmax_gemm_permute/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_f16_f16_f16_f16_gmk_gnk_gno_gmo_instance.cpp
View file @
5f4a0f73
...
...
@@ -65,8 +65,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskOutUpperTriangle
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
void
add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_instances
(
std
::
vector
<
...
...
@@ -106,8 +106,8 @@ void add_device_batched_gemm_softmax_gemm_permute_xdl_cshuffle_gmk_gnk_gno_gmo_i
PassThrough
,
PassThrough
,
MaskingSpecialization
::
MaskDisabled
>
;
DeviceOperationInstanceCreator
<
DeviceOp
,
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
>::
add_device_instances
(
instances
);
DeviceOperationInstanceCreator
<
ArchitectureEnumSequence
<
ArchitectureEnum
::
Gfx908
>
,
DeviceOp
>::
add_device_instances
(
instances
);
}
}
// namespace instance
...
...
src_example/01_gemm_bias_softmax_gemm_permute/gemm_bias_softmax_gemm_permute.cpp
View file @
5f4a0f73
...
...
@@ -136,9 +136,9 @@ int main()
// get device op instances
std
::
vector
<
std
::
unique_ptr
<
DeviceOp
>>
op_ptrs
;
ck
::
tensor_operation
::
device
::
instance
::
DeviceOperationInstanceCreator
<
DeviceOp
,
ck
::
tensor_operation
::
device
::
instance
::
GemmFea
tureEnum
::
Xdl
>::
add_device_instances
(
op_ptrs
);
ck
::
tensor_operation
::
device
::
instance
::
ArchitectureEnumSequence
<
ck
::
tensor_operation
::
device
::
instance
::
Architec
tureEnum
::
All
>
,
DeviceOp
>::
add_device_instances
(
op_ptrs
);
std
::
cout
<<
"found "
<<
op_ptrs
.
size
()
<<
" instances"
<<
std
::
endl
;
...
...
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