supporting different hidden dimensions (microsoft#559)

* supporting different hidden dimensions * add support for larger hidden dimensions (greater than 8K) * remove empty line * add loop unrolling factor for dropout kernels * update different kernels based on the reviews Co-authored-by: Jeff Rasley <[email protected]>
bobisai · Dec 1, 2020 · c78c29f · c78c29f
1 parent 17f36f1
commit c78c29f
Show file tree

Hide file tree

Showing 8 changed files with 1,027 additions and 654 deletions.
diff --git a/csrc/includes/custom_cuda_layers.h b/csrc/includes/custom_cuda_layers.h
@@ -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,

diff --git a/csrc/transformer/cublas_wrappers.cu b/csrc/transformer/cublas_wrappers.cu
@@ -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,