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
4b8f1249
"git@developer.sourcefind.cn:OpenDAS/mmcv.git" did not exist on "3f8e985b02d7ef1c0b3a3461e872883a25027bee"
Commit
4b8f1249
authored
Dec 16, 2022
by
fsx950223
Browse files
fix format
parent
19c18624
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
31 additions
and
179 deletions
+31
-179
example/36_sparse_embedding/sparse_embedding3_forward_layernorm.cpp
..._sparse_embedding/sparse_embedding3_forward_layernorm.cpp
+31
-179
No files found.
example/36_sparse_embedding/sparse_embedding3_forward_layernorm.cpp
View file @
4b8f1249
...
@@ -25,172 +25,25 @@ using BetaDataType = ck::half_t;
...
@@ -25,172 +25,25 @@ using BetaDataType = ck::half_t;
using
AccDataType
=
float
;
using
AccDataType
=
float
;
using
OutType
=
ck
::
half_t
;
using
OutType
=
ck
::
half_t
;
using
DeviceInstance_fp16_e256
=
using
DeviceInstance_fp16_e256
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
256
,
1
,
1
,
3
>
;
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
using
DeviceInstance_fp16_e512
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
512
,
1
,
2
,
3
>
;
IndexType
,
using
DeviceInstance_fp16_e768
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
768
,
1
,
1
,
3
>
;
GammaDataType
,
using
DeviceInstance_fp16_e1024
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
1024
,
1
,
2
,
3
>
;
BetaDataType
,
using
DeviceInstance_fp16_e1536
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
1536
,
1
,
2
,
3
>
;
AccDataType
,
using
DeviceInstance_fp16_e2048
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
2048
,
1
,
2
,
3
>
;
OutType
,
using
DeviceInstance_fp16_e4096
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
4096
,
1
,
8
,
3
>
;
256
,
using
DeviceInstance_fp16_e8192
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
8192
,
1
,
8
,
3
>
;
1
,
256
,
template
<
typename
emb_type
,
ck
::
index_t
dim
>
struct
emb_kernel
{};
1
,
256
,
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
256
>
{
using
kernel_type
=
DeviceInstance_fp16_e256
;
};
1
,
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
512
>
{
using
kernel_type
=
DeviceInstance_fp16_e512
;
};
1
,
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
768
>
{
using
kernel_type
=
DeviceInstance_fp16_e768
;
};
3
>
;
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
1024
>
{
using
kernel_type
=
DeviceInstance_fp16_e1024
;
};
using
DeviceInstance_fp16_e512
=
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
1536
>
{
using
kernel_type
=
DeviceInstance_fp16_e1536
;
};
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
2048
>
{
using
kernel_type
=
DeviceInstance_fp16_e2048
;
};
IndexType
,
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
4096
>
{
using
kernel_type
=
DeviceInstance_fp16_e4096
;
};
GammaDataType
,
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
8192
>
{
using
kernel_type
=
DeviceInstance_fp16_e8192
;
};
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
512
,
1
,
2
,
3
>
;
using
DeviceInstance_fp16_e768
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
768
,
1
,
1
,
3
>
;
using
DeviceInstance_fp16_e1024
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
1024
,
1
,
2
,
3
>
;
using
DeviceInstance_fp16_e1536
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
1536
,
1
,
2
,
3
>
;
using
DeviceInstance_fp16_e2048
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
2048
,
1
,
2
,
3
>
;
using
DeviceInstance_fp16_e4096
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
4096
,
1
,
8
,
3
>
;
using
DeviceInstance_fp16_e8192
=
ck
::
tensor_operation
::
device
::
DeviceSparseEmbeddingsForwardLayernorm
<
EmbType
,
IndexType
,
GammaDataType
,
BetaDataType
,
AccDataType
,
OutType
,
256
,
1
,
256
,
1
,
8192
,
1
,
8
,
3
>
;
template
<
typename
emb_type
,
ck
::
index_t
dim
>
struct
emb_kernel
{
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
256
>
{
using
kernel_type
=
DeviceInstance_fp16_e256
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
512
>
{
using
kernel_type
=
DeviceInstance_fp16_e512
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
768
>
{
using
kernel_type
=
DeviceInstance_fp16_e768
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
1024
>
{
using
kernel_type
=
DeviceInstance_fp16_e1024
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
1536
>
{
using
kernel_type
=
DeviceInstance_fp16_e1536
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
2048
>
{
using
kernel_type
=
DeviceInstance_fp16_e2048
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
4096
>
{
using
kernel_type
=
DeviceInstance_fp16_e4096
;
};
template
<
>
struct
emb_kernel
<
ck
::
half_t
,
8192
>
{
using
kernel_type
=
DeviceInstance_fp16_e8192
;
};
// clang-format on
// clang-format on
...
@@ -270,19 +123,18 @@ int main()
...
@@ -270,19 +123,18 @@ int main()
beta_dev
.
ToDevice
(
beta
.
mData
.
data
());
beta_dev
.
ToDevice
(
beta
.
mData
.
data
());
auto
device_instance
=
typename
emb_kernel
<
EmbType
,
current_dim
>::
kernel_type
{};
auto
device_instance
=
typename
emb_kernel
<
EmbType
,
current_dim
>::
kernel_type
{};
auto
argument_ptr
=
device_instance
.
MakeArgumentPointer
(
auto
argument_ptr
=
device_instance
.
MakeArgumentPointer
(
out_dev
.
GetDeviceBuffer
(),
out_dev
.
GetDeviceBuffer
(),
{
ck
::
type_convert
<
EmbType
*>
(
emb_a_dev
.
GetDeviceBuffer
()),
{
ck
::
type_convert
<
EmbType
*>
(
emb_a_dev
.
GetDeviceBuffer
()),
ck
::
type_convert
<
EmbType
*>
(
emb_b_dev
.
GetDeviceBuffer
()),
ck
::
type_convert
<
EmbType
*>
(
emb_b_dev
.
GetDeviceBuffer
()),
ck
::
type_convert
<
EmbType
*>
(
emb_c_dev
.
GetDeviceBuffer
())},
ck
::
type_convert
<
EmbType
*>
(
emb_c_dev
.
GetDeviceBuffer
())},
{
ck
::
type_convert
<
IndexType
*>
(
index_a_dev
.
GetDeviceBuffer
()),
{
ck
::
type_convert
<
IndexType
*>
(
index_a_dev
.
GetDeviceBuffer
()),
ck
::
type_convert
<
IndexType
*>
(
index_b_dev
.
GetDeviceBuffer
()),
ck
::
type_convert
<
IndexType
*>
(
index_b_dev
.
GetDeviceBuffer
()),
ck
::
type_convert
<
IndexType
*>
(
index_c_dev
.
GetDeviceBuffer
())},
ck
::
type_convert
<
IndexType
*>
(
index_c_dev
.
GetDeviceBuffer
())},
gamma_dev
.
GetDeviceBuffer
(),
gamma_dev
.
GetDeviceBuffer
(),
beta_dev
.
GetDeviceBuffer
(),
beta_dev
.
GetDeviceBuffer
(),
current_dim
,
current_dim
,
index_length
,
index_length
,
epsilon
);
epsilon
);
std
::
cout
<<
"Dim:"
<<
current_dim
<<
", kernel:"
<<
device_instance
.
GetTypeString
()
std
::
cout
<<
"Dim:"
<<
current_dim
<<
", kernel:"
<<
device_instance
.
GetTypeString
()
<<
std
::
endl
<<
std
::
endl
<<
std
::
flush
;
<<
std
::
flush
;
...
...
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