-
Notifications
You must be signed in to change notification settings - Fork 666
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Positional encoding #9772
base: master
Are you sure you want to change the base?
Positional encoding #9772
Conversation
编写fused_sinusoidal_positional_encode实现如下逻辑 |
); | ||
let attrs = (ins | ||
SI32Attr:$embedding_dim, | ||
DefaultValuedAttr<SI32Attr, "0">:$pattern, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
pattern是不是换成layout更好
这里可以考虑传字符串
int max_period; | ||
}; | ||
|
||
enum class EncodingPattern { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
enum class EncodingPattern { | |
enum class EncodingLayout{ |
}; | ||
|
||
enum class EncodingPattern { | ||
SIN_COS, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
SIN_COS, | |
kPlanarSinCos, |
} | ||
} | ||
|
||
__global__ void PaddingKernel(float* out_ptr, int N, int embedding_dim) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
避免用大写的 N
DispatchSrcType(positions->data_type(), EncodeType::SIN, sin_param); | ||
DispatchSrcType(positions->data_type(), EncodeType::COS, cos_param); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
这里调用两次,就会有两次的启动开销,同时是不是还会有非连续访存的问题,尽量使用一次kernel launch,同时保证连续访存
No description provided.