Skip to content
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

Draft
wants to merge 8 commits into
base: master
Choose a base branch
from
Draft

Positional encoding #9772

wants to merge 8 commits into from

Conversation

yz-chen18
Copy link
Contributor

No description provided.

@yz-chen18
Copy link
Contributor Author

yz-chen18 commented Jan 18, 2023

编写fused_sinusoidal_positional_encode实现如下逻辑
https://github.com/Oneflow-Inc/diffusers/blob/oneflow-fork/src/diffusers/models/embeddings_oneflow.py#L21
包含kPlanarSinCos, kPlanarCosSin, kInterleavedSinCos, kInterleavedCosSin四种layout
幻灯片1
幻灯片2

);
let attrs = (ins
SI32Attr:$embedding_dim,
DefaultValuedAttr<SI32Attr, "0">:$pattern,
Copy link
Collaborator

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 {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
enum class EncodingPattern {
enum class EncodingLayout{

};

enum class EncodingPattern {
SIN_COS,
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
SIN_COS,
kPlanarSinCos,

}
}

__global__ void PaddingKernel(float* out_ptr, int N, int embedding_dim) {
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

避免用大写的 N

Comment on lines 146 to 147
DispatchSrcType(positions->data_type(), EncodeType::SIN, sin_param);
DispatchSrcType(positions->data_type(), EncodeType::COS, cos_param);
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这里调用两次,就会有两次的启动开销,同时是不是还会有非连续访存的问题,尽量使用一次kernel launch,同时保证连续访存

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants