Skip to content

supporting different hidden dimensions #559

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.

Sign up for GitHub

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

Merged
RezaYazdaniAminabadi merged 9 commits into from
Dec 1, 2020
Merged

supporting different hidden dimensions #559

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
de20c4c
supporting different hidden dimensions
Nov 25, 2020
98764b9
Merge branch 'master' into transformer-kernel/arbitrary-hidden-dimension
jeffra Nov 25, 2020
0443a55
Merge branch 'master' into transformer-kernel/arbitrary-hidden-dimension
jeffra Nov 25, 2020
733927d
add support for larger hidden dimensions (greater than 8K)
Nov 25, 2020
539f2e3
Merge branch 'transformer-kernel/arbitrary-hidden-dimension' of githu…
Nov 25, 2020
86c7735
remove empty line
Nov 25, 2020
d50fbb4
add loop unrolling factor for dropout kernels
Nov 26, 2020
f3dd328
update different kernels based on the reviews
Dec 1, 2020
21dd505
Merge branch 'master' into transformer-kernel/arbitrary-hidden-dimension
jeffra Dec 1, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions csrc/includes/custom_cuda_layers.h
100644 → 100755
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,8 @@
#define MAX_THREAD_ITERATIONS 8 // Maximum 8K
#define MAX_WARP_NUM 32

#define MAX_REGISTERS 256

// Fused bias add with gelu activation
template <typename T>
void launch_bias_gelu(const T* input,
Expand Down
3 changes: 2 additions & 1 deletion csrc/transformer/cublas_wrappers.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -133,7 +133,8 @@ int cublas_strided_batched_gemm(cublasHandle_t handle,

if (status != CUBLAS_STATUS_SUCCESS) {
fprintf(stderr,
"!!!! kernel execution error. (m: %d, n: %d, k: %d, error: %d) \n",
"!!!! kernel execution error. (batch: %d, m: %d, n: %d, k: %d, error: %d) \n",
batch,
m,
n,
k,
Expand Down
Loading