comp files

comp.sh

@@ -0,0 +1,14 @@
+cp kernels_req_1/attention.cuh kernels/attention.cuh 
+cp kernels_req_1/matmul.cuh kernels/matmul.cuh 
+cp kernels_req_3/layernorm.cuh kernels/layernorm.cuh
+cp kernels_req_3/softmax.cuh kernels/softmax.cuh
+
+make test_gpt2
+
+git checkout kernels/attention.cuh
+git checkout kernels/matmul.cuh
+
+echo "RUNNING !!!"
+sbatch verify.slurm
+
+squeue -u $USER

gpt2.cuh

@@ -150,22 +150,23 @@ void fill_in_activation_sizes(size_t* act_sizes, int B, int T, GPT2Config config
     size_t C = config.channels;
     act_sizes[0] = B * T * C; // encoded
     act_sizes[1] = L * B * T * C; // ln1
-    act_sizes[2] = 1; // ln1_mean, reserved for future use
+    //to align making 1 to 4
-    act_sizes[3] = 1; // ln1_rstd, reserved for future use
+    act_sizes[2] = 4; // ln1_mean, reserved for future use
+    act_sizes[3] = 4; // ln1_rstd, reserved for future use
     act_sizes[4] = L * B * T * C; // atty
     act_sizes[5] = L * B * NH * T * T; // att
     act_sizes[6] = L * B * T * C; // attproj
     act_sizes[7] = L * B * T * C; // residual2
     act_sizes[8] = L * B * T * C; // ln2
-    act_sizes[9] = 1; // ln2_mean, reserved for future use
+    act_sizes[9] = 4; // ln2_mean, reserved for future use
-    act_sizes[10] = 1; // ln2_rstd, reserved for future use
+    act_sizes[10] = 4; // ln2_rstd, reserved for future use
     act_sizes[11] = L * B * T * 4*C; // fch
     act_sizes[12] = L * B * T * 4*C; // fch_gelu
     act_sizes[13] = L * B * T * C; // fcproj
     act_sizes[14] = L * B * T * C; // residual3
     act_sizes[15] = B * T * C; // lnf
-    act_sizes[16] = 1; // lnf_mean, reserved for future use
+    act_sizes[16] = 4; // lnf_mean, reserved for future use
-    act_sizes[17] = 1; // lnf_rstd, reserved for future use
+    act_sizes[17] = 4; // lnf_rstd, reserved for future use
     act_sizes[18] = B * T; // losses
     act_sizes[19] = L * B * T * 3*C; // qkvr
     act_sizes[20] = B * T * max(3*C, max(NH*T, Vp)); // output / scratch
@@ -179,19 +180,23 @@ typedef struct {
     float* residual3; // (B, T, C)
 } GradActTensors;
 
+// op_13: 256-byte aligned activations.
+// Each activation tensor's stride is rounded up to a multiple of 64 floats (256 bytes).
+#define align_to_256_bytes(x) \
+    (x + (256/sizeof(float) - (x % (256/sizeof(float)))) % (256/sizeof(float)))
+
 float* malloc_and_point(float** targets[], const size_t* act_sizes, int n) {
     size_t num_activations = 0;
     for (size_t i = 0; i < n; i++) {
-        num_activations += act_sizes[i];
+        num_activations += align_to_256_bytes(act_sizes[i]);
     }
     float* acts_memory;
     cudaCheck(cudaMalloc((void**)&acts_memory, num_activations * sizeof(float)));
-    // clear the allocated memory since cudaMalloc does not guarantee zeroed memory
     cudaCheck(cudaMemset(acts_memory, 0, num_activations * sizeof(float)));
     float* acts_memory_iterator = acts_memory;
     for (size_t i = 0; i < n; i++) {
         *(targets[i]) = acts_memory_iterator;
-        acts_memory_iterator += act_sizes[i];
+        acts_memory_iterator += align_to_256_bytes(act_sizes[i]);
     }
     return acts_memory;
 }
@@ -385,8 +390,10 @@ void gpt2_forward(GPT2 *model, int* inputs, int B, int T) {
         matmul_forward(l_attproj, l_atty, l_attprojw, l_attprojb, B, T, C, C);
         residual_forward(l_residual2, residual, l_attproj, B*T*C);
         layernorm_forward(l_ln2, l_ln2_mean, l_ln2_rstd, l_residual2, l_ln2w, l_ln2b, B, T, C);
-        matmul_forward(l_fch, l_ln2, l_fcw, l_fcb, B, T, C, 4*C);
+        //op16 matmul + bias + GeLU fusion
-        gelu_forward(l_fch_gelu, l_fch, B*T*4*C);
+        // matmul_forward(l_fch, l_ln2, l_fcw, l_fcb, B, T, C, 4*C);
+        // gelu_forward(l_fch_gelu, l_fch, B*T*4*C);
+        matmul_forward_fused_gelu(l_fch_gelu, l_ln2, l_fcw, l_fcb, B, T, C, 4*C);
         matmul_forward(l_fcproj, l_fch_gelu, l_fcprojw, l_fcprojb, B, T, 4*C, C);
         residual_forward(l_residual3, l_residual2, l_fcproj, B*T*C);
     }

kernels/attention.cuh

@@ -2,10 +2,159 @@
 #define __ATTENTION_CUH__
 
 #include <cuda_runtime.h>
+#include <math.h>
+#include <float.h>
 #include "../utils/cuda_utils.cuh"
-#include "softmax.cuh"
 
-__global__ void permute_kernel(float* q, float* k, float* v, const float* inp, int B, int N, int NH, int d) {
+
+//added templates so the compiler can unroll loops and optimize better. 
+// op11
+// without swilling the QK^T inner loops read smem_q [query_row][i] where each thread query_row in teh warp holds a different row but reads teh same column i. as HS is 64 every thread bank index reduces to (i % 32) giving 32 way back conflict per warp on the Q reads. col -> col ^ row to swizzle the accesses and reduce bank conflicts. so each 32 read to 32 different banks instead of 32 reads to the same bank.
+template <int BR, int BC, int HS>
+__global__ void flash_attention_forward_kernel(const float* __restrict__ Q, const float* __restrict__ K, const float* __restrict__ V, float* __restrict__ out, int B, int T, int NH, float scale) {
+    //following line by like psudo code from flash attention paper.
+    // https://tridao.me/publications/flash2/flash2.pdf
+
+    // step 1 map thread blocks to (batch , head , query_tile)
+    int b = blockIdx.x;
+    int nh = blockIdx.y;
+    int query_tile = blockIdx.z; 
+
+    int query_row = threadIdx.x; //load index
+    int query_global = query_tile * BR + query_row; 
+
+    //step 2 calculate all the base pointers for this block. We will be loading queries from Q and keys/values from K and V, and writing output to out.
+    const float* __restrict__ query_base = Q + b * NH * T * HS + nh * T * HS;
+    const float* __restrict__ key_base = K + b * NH * T * HS + nh * T * HS;
+    const float* __restrict__ value_base = V + b * NH * T * HS + nh * T * HS;
+    float* out_base = out + b * NH * T * HS + nh * T * HS;
+
+    // declare shared memory for the tile of queries and keys/values. We will load a tile of queries and then iterate over tiles of keys/values.
+    __shared__ float smem_q[BR * HS];
+    __shared__ float smem_k[BC * HS];
+    __shared__ float smem_v[BC * HS];
+
+    // we will compute the output for a tile of queries at a time.
+    int query_tile_start = query_tile * BR; // global index of the first query in this tile. This is used to load the correct queries into shared memory and also to know which output rows we are responsible for computing.
+    bool q_valid = (query_global < T); // check if this thread's query is within the sequence length. If not, it will do nothing.
+
+
+    //define swizzling function using the XOR on col and row index to reduce bank conflicts on smem_q reads in the inner loop.
+
+    #define SWIZZLE(col, row) ((row) * HS + ((col) ^ ((row) & (HS - 1))))
+
+    // step 4 load query tiles from global memory to shared memory.
+    // coalesced load of the queries for this tile into shared memory. a linear load of BR*HS elements, with each thread loading HS elements for one query.
+    #pragma unroll
+    for (int i = 0; i < HS; i++) {
+        smem_q[SWIZZLE(i, query_row)] = (query_tile_start + query_row < T) ? query_base[(query_tile_start + query_row) * HS + i] : 0.0f; // if the query is out of bounds, we can load zeros since they won't contribute to the output.
+    }
+
+    // step 5 init accumulators for the output of this tile of queries.
+    float out_reg[HS];
+    #pragma unroll  
+    for (int i = 0; i < HS; i++) {
+        out_reg[i] = 0.0f;
+    }
+    float max_score = -FLT_MAX;
+    float sum_exp = 0.0f;
+
+
+    // step 6 determine KV tile range for casual masking
+    int query_tile_end = (query_tile + 1) * BR;
+    int n_kv_tiles = (query_tile_end < T) ? (T + BC - 1) / BC : (query_tile_end + BC - 1) / BC; // if the end of the query tile is within the sequence length, we only need to iterate over the KV tiles that cover the sequence length. If the end of the query tile exceeds the sequence length, we need to iterate over all KV tiles that cover the end of the query tile to properly apply casual masking.
+
+    __syncthreads();
+
+    //step 7 main loop over kv values 
+    for(int j = 0; j < n_kv_tiles; j++){
+        __syncthreads();
+
+        int k_tile_start = j * BC;
+
+        // step 8 load k and v tiles from global memory
+        #pragma unroll
+        for (int i = 0; i < HS; i++) {
+            int l = i * BC + query_row;
+            int row = l / HS;
+            int col = l % HS;
+            int global_row = k_tile_start + row;
+            bool valid = (global_row < T) && (row < BC);
+
+            if (valid) {
+                smem_k[SWIZZLE(col, row)] = key_base[global_row * HS + col];
+                smem_v[SWIZZLE(col, row)] = value_base[global_row * HS + col];
+            } else {
+                smem_k[SWIZZLE(col, row)] = 0.0f; // pad with zeros if out of bounds
+                smem_v[SWIZZLE(col, row)] = 0.0f; // pad with zeros if out of bounds
+            }
+        }
+        __syncthreads();
+
+        // step 9 compute one row of S = Qi * Kj^T and apply casual masking
+        float score_reg[BC];
+        float row_max = -FLT_MAX;
+
+        #pragma unroll
+        for (int c = 0; c < BC; c++) {
+            int k_global = j * BC + c;
+            float dot; // compute the dot product of the query and key for this position
+
+            if (!q_valid || k_global > query_global) { // apply casual masking. If the key position is greater than the query position, we mask it out by setting the score to a very large negative number.
+                dot = -FLT_MAX;
+            } else {
+                float acc = 0.0f;
+                #pragma unroll
+                for (int h = 0; h < HS; h++) {
+                    acc += smem_q[SWIZZLE(h, query_row)] * smem_k[SWIZZLE(h, c)];
+                }
+                dot = acc * scale;
+            }
+            score_reg[c] = dot;
+            row_max = fmaxf(row_max, dot);  // local max for this row
+        }
+
+        // step 10 online softmax rescaling 
+        float max_new = fmaxf(max_score, row_max);
+        if (max_new == -FLT_MAX) { // if both are -FLT_MAX, we can skip the rescaling because all scores are -FLT_MAX and the output will be zero anyway.
+            continue; 
+        }
+        float alpha = (max_score == -FLT_MAX) ? 0.0f : expf(max_score - max_new); // if max_score is -FLT_MAX, we set alpha to 0 to avoid exp overflow. Otherwise, we compute the scaling factor for the old scores.
+        float local_tile = 0.0f; // this will hold the sum of exp for the scores in this tile after rescaling.
+
+        #pragma unroll
+        for (int c = 0; c < BC; c++) {
+            float p = (score_reg[c] == -FLT_MAX) ? 0.0f : expf(score_reg[c] - max_new); // compute the exp for this score after rescaling. If the score is -FLT_MAX, we set it to 0 to avoid exp underflow.
+            local_tile += p;
+            score_reg[c] = p; // store the rescaled exp in score_reg for later use in the weighted sum.
+        }
+
+        // update output acc 
+        #pragma unroll
+        for (int h = 0; h < HS; ++h) {
+            float pv = 0.0f;
+            #pragma unroll
+            for (int c = 0; c < BC; ++c) {
+                pv += score_reg[c] * smem_v[SWIZZLE(h, c)];
+            }
+            out_reg[h] = alpha * out_reg[h] + pv;
+        }
+        sum_exp = alpha * sum_exp + local_tile;  // update the local tile sum with the contribution from this tile and the previous tiles, scaled by alpha.
+        max_score = max_new; // update the max score for the next iteration
+
+    }
+    if (q_valid) {
+        float inv_l = (sum_exp > 0.0f) ? (1.0f / sum_exp) : 0.0f;
+        #pragma unroll
+        for (int h = 0; h < HS; ++h) {
+            out_base[query_global * HS + h] = out_reg[h] * inv_l;
+        }
+    }
+}
+
+
+
+__global__ void permute_kernel(float* __restrict__ q, float* __restrict__ k, float*  __restrict__ v, const float* __restrict__ inp, int B, int N, int NH, int d) {
     // Implement this
 
     int x = blockDim.x * blockIdx.x + threadIdx.x;
@@ -24,7 +173,7 @@ __global__ void permute_kernel(float* q, float* k, float* v, const float* inp, i
 }
 
 
-__global__ void unpermute_kernel(float* inp, float *out, int B, int N, int NH, int d) {
+__global__ void unpermute_kernel(const float* __restrict__ inp, float* __restrict__ out, int B, int N, int NH, int d) {
     // Implement this
     int x = blockDim.x * blockIdx.x + threadIdx.x;
 
@@ -39,39 +188,42 @@ __global__ void unpermute_kernel(float* inp, float *out, int B, int N, int NH, i
 
 }
 
-__global__ void preattention_scores_kernel(float* q, float* k, int B, int T, int NH, int HS, float* preatt){
+// __global__ void preattention_scores_kernel(const float* __restrict__ q, const float* __restrict__ k, int B, int T, int NH, int HS, float* __restrict__ preatt){
-    int x = blockDim.x * blockIdx.x + threadIdx.x;
+//     int x = blockDim.x * blockIdx.x + threadIdx.x;
 
-    if(x < B* T* NH* T){
+//     if(x < B* T* NH* T){
-        int b = (x / T / T / NH);
+//         int b = (x / T / T / NH);
-        int nh = (x / T / T) % NH;
+//         int nh = (x / T / T) % NH;
-        int t1 = (x / T) % T;
+//         int t1 = (x / T) % T;
-        int t2 = x % T;
+//         int t2 = x % T;
 
-        float sum = 0.0f;
+//         float sum = 0.0f;
-        for (int i = 0; i < HS; i++) {
+//         const float* __restrict__ q_ptr = q + b * NH * T * HS + nh * T * HS + t1 * HS;
-            sum += q[b*NH*T*HS + nh*T*HS + t1*HS + i] * k[b*NH*T*HS + nh*T*HS + t2*HS + i];
+//         const float* __restrict__ k_ptr = k + b * NH * T * HS + nh * T * HS + t2 * HS;
-        }
+//         for (int i = 0; i < HS; i++) {
-        preatt[b * NH * T * T + nh * T * T + t1 * T + t2] = sum;
+//             sum += q_ptr[i] * k_ptr[i];
-    }   
+//         }
-}
+//         preatt[b * NH * T * T + nh * T * T + t1 * T + t2] = sum;
+//     }   
+// }
 
-__global__ void attention_v_kernel(float* vaccum, float* att, float* v, int B, int T, int NH, int HS){
+// __global__ void attention_v_kernel(float* __restrict__ vaccum, const float* __restrict__ att, const float* __restrict__ v, int B, int T, int NH, int HS){
-    int x = blockDim.x * blockIdx.x + threadIdx.x;
+//     int x = blockDim.x * blockIdx.x + threadIdx.x;
 
-    if(x < B*NH*T*HS){
+//     if(x < B*NH*T*HS){
-        int b   = (x / HS / T / NH);
+//         int b   = (x / HS / T / NH);
-        int nh  = (x / HS / T) % NH;
+//         int nh  = (x / HS / T) % NH;
-        int t   = (x / HS) % T;
+//         int t   = (x / HS) % T;
-        int hs  = x % HS;
+//         int hs  = x % HS;
-        float sum = 0.0f;
+//         float sum = 0.0f;
-        for (int t2 = 0; t2 < T; t2++) {
+//         const float* __restrict__ att_ptr = att + b * NH * T * T + nh * T * T + t * T;
-            sum += att[b*NH*T*T + nh*T*T + t*T + t2] * v[b*NH*T*HS + nh*T*HS + t2*HS + hs];
+//         for (int t2 = 0; t2 < T; t2++) {
-        }
+//             sum += att_ptr[t2] * v[b*NH*T*HS + nh*T*HS + t2*HS + hs];
+//         }
 
-        vaccum[b * NH * T * HS + nh * T * HS + t * HS + hs] = sum;
+//         vaccum[b * NH * T * HS + nh * T * HS + t * HS + hs] = sum;
-    }
+//     }
-}
+// }
 
 // Launch all kernels related to attention here 
 void attention_forward(float* out, float* qkvr, float* att, float* inp, int B, int T, int C, int NH) {
@@ -83,21 +235,23 @@ void attention_forward(float* out, float* qkvr, float* att, float* inp, int B, i
     k = qkvr + 1 * B * T * C;
     v = qkvr + 2 * B * T * C;
 
-    dim3 block(32,1,1);
+    dim3 block(256,1,1);
     dim3 grid((B * NH * T * HS + block.x-1) / block.x,1,1);
 
     permute_kernel<<<grid, block>>>(q, k, v, inp, B, T, NH, HS);
 
-    float* preatt = inp;
+    //fa2 — op_7 sweep winner: BR=32, BC=32 was 26% faster than BR=64,BC=64
-    dim3 grid_preatt((B * NH * T * T + block.x-1) / block.x,1,1);
+    // at T=64 (smaller tiles fit two blocks per SM, doubling occupancy).
-    preattention_scores_kernel<<<grid_preatt, block>>>(q, k, B, T, NH, HS, preatt);
+    constexpr int BR = 32; // tile size for queries
+    constexpr int BC = 32; // tile size for keys/values
+    float scale = 1.0f / sqrtf((float)HS);
 
-    float scale = 1.0 / sqrtf(HS);
+    int number_of_q_tiles = (T + BR - 1) / BR;
-    dim3 grid_softmax (( B * NH * T + block.x -1 )/ block.x,1,1);
+    dim3 fa_grid(B, NH, number_of_q_tiles);      // one block per (batch, head, Q tile)
-    softmax_forward_kernel<<<grid_softmax, block>>>(att, scale, preatt, B * NH, T);
+    dim3 fa_block(BR, 1, 1); // one warp per query in the tile
 
     float* vaccum = inp;
-    attention_v_kernel<<<grid_preatt,block>>>(vaccum, att, v, B, T, NH, HS);
+    flash_attention_forward_kernel<BR, BC, 64><<<fa_grid, fa_block>>>(q, k, v, vaccum, B, T, NH, scale);
 
     unpermute_kernel<<<grid, block>>>(vaccum, out, B, T, NH, HS);
 }

kernels/attention_kv_cache_template.cuh

@@ -1,86 +0,0 @@
-#ifndef __ATTENTION_KV_CACHE_TEMPLATE_CUH__
-#define __ATTENTION_KV_CACHE_TEMPLATE_CUH__
-
-#include <cuda_runtime.h>
-#include "../utils/cuda_utils.cuh"
-
-extern bool g_enable_kv_cache; // Toggles KV-caching logic
-extern float* g_kv_cache;      // Pointer to the GPU-allocated KV cache buffer
-extern int g_layer_idx;        // Tracks the current transformer layer, useful for knowing when to switch to decode
-extern bool g_is_prefill;      // True during the initial prompt, False during token generation
-extern int g_current_pos;      // Tracks the absolute sequence position during Decode, helps index into the KV cache
-
-// --- Suggested Kernel Placeholders (TODO: Add your regular attention kernel here, and any other kernels you may need) ---
-
-/**
- * Kernel to store the newly computed Key and Value into the global KV cache. 
- * CAN BE FUSED WITH QKV MATMUL
- * In Prefill: Stores T tokens for the current layer.
- * In Decode: Stores 1 token at g_current_pos for the current layer.
- */
-__global__ void store_kv_kernel(/* args */) {
-    // Implement logic to copy K and V from the input QKV buffer into g_kv_cache
-}
-
-/**
- * Kernel to compute the attention for the Decode phase (T=1).
- * This involves:
- * 1. Extracting the current Query (Q)
- * 2. Gathering all previous Keys (K) and Values (V) from the cache
- * 3. Computing attention using the new Q and cached KV values.
- */
-__global__ void decode_attention_kernel(/* args */) {
-    // Implement T=1 attention logic using the KV cache
-}
-
-void attention_forward(float* out, float* qkvr, float* att, float* inp, int B, int T, int C, int NH) {
-    // 1. Check if KV-caching is enabled, used by the verification script
-    if (!g_enable_kv_cache) {
-        // TODO: Call your standard Baseline Attention here (Milestone 1/2)
-        return;
-    }
-
-    // 2. Initialize KV Cache if necessary
-    if (g_kv_cache == nullptr) {
-        // TODO: Allocate the KV cache, based accordingly to the model parameters
-        // Option 1: Allocate a very big kv cache to store all the tokens
-        // Option 2: Dynamically resize the kv cache repeatedly for each new token
-    }
-
-    if (g_is_prefill) {
-        // --- PREFILL PHASE ---
-        // During prefill, we process the entire input sequence (T > 1).
-        
-        // A. Store the K and V results from this layer into the cache.
-        //    The projections are available in 'inp' (the QKV matmul output).
-        
-        // B. Run standard baseline attention for the initial prompt tokens.
-
-        // C. Update the global layer index to track progress through the model.
-        g_layer_idx++;
-        if (g_layer_idx == 12) {
-            g_layer_idx = 0;
-            g_is_prefill = false; // Transition to Decode phase for the next token
-        }
-    } else {
-        // --- DECODE PHASE ---
-        // During decode, we only process the single newest token (T = 1).
-        
-        // A. Store the single new Key and Value for this step into the cache at g_current_pos.
-        
-        // B. Run specialized Decode Attention.
-        //    This requires reading the history of Keys and Values from the cache
-        //    to compute the output for the current token.
-        //    TODO: Don't forget to scale your attention scores
-        //    TODO: Ensure the output shapes match the next layer
-        
-        // C. Update the global state.
-        g_layer_idx++;
-        if (g_layer_idx == 12) {
-            g_layer_idx = 0;
-            g_current_pos++; // Advance the sequence position
-        }
-    }
-}
-
-#endif // __ATTENTION_KV_CACHE_TEMPLATE_CUH__

kernels/encoder.cuh

@@ -4,27 +4,31 @@
 #include <cuda_runtime.h>
 #include "../utils/cuda_utils.cuh"
 
-__global__ void encoder_forward_kernel(float* out, const int* inp, const float* wte, const float* wpe,
+__global__ void encoder_forward_kernel(float* __restrict__ out, const int* __restrict__ inp, const float* __restrict__ wte, const float* __restrict__ wpe,
                                        int B, int T, int C) {
     // Implement this
-    int x = blockDim.x * blockIdx.x + threadIdx.x;
+    int row = blockIdx.x * blockDim.y + threadIdx.y;
-    int y = blockDim.y * blockIdx.y + threadIdx.y;
+    if (row >= B * T){
-
+        return;
-    if(x < B && y < T){
+    }
-        int token = inp[x*T + y];
+    int b = row / T;
-        for(int i = 0; i < C; i ++){
+    int t = row % T;
-            out[x*T*C + y*C + i] = wte[token*C+i] + wpe[y*C+i];
+    int token = inp[b * T + t];
-        }
+    const float* __restrict__ wte_ptr = wte + token * C;
+    const float* __restrict__ wpe_ptr = wpe + t * C;
+    float* __restrict__ out_ptr = out + row * C;
+    for (int i = threadIdx.x; i < C; i += 32) {
+        out_ptr[i] = wte_ptr[i] + wpe_ptr[i];
     }
 }
 
 // Launch kernel here
 void encoder_forward(float* out, const int* inp, const float* wte, const float* wpe, int B, int T, int C) {
     // Implement this
-    dim3 block(16,16,1);
+    dim3 block(32, 8, 1);
-    dim3 grid((B + block.x -1)/ block.x,(T+block.y-1)/block.y);
+    int total_rows = B * T;
-
+    dim3 grid((total_rows + 8 - 1) / 8, 1, 1);
-    encoder_forward_kernel<<<grid, block>>>(out, inp, wte, wpe,B,T,C);
+    encoder_forward_kernel<<<grid, block>>>(out, inp, wte, wpe, B, T, C);
 }
 
 

kernels/gelu.cuh

@@ -6,7 +6,7 @@
 
 #define GELU_SCALING_FACTOR sqrtf(2.0f / M_PI)
 
-__global__ void gelu_forward_kernel(float* out, const float* inp, int N) {
+__global__ void gelu_forward_kernel(float* __restrict__ out, const float* __restrict__ inp, int N) {
     // Implement this
 
     int x = blockDim.x * blockIdx.x + threadIdx.x;
@@ -22,7 +22,7 @@ __global__ void gelu_forward_kernel(float* out, const float* inp, int N) {
 // Launch kernel here
 void gelu_forward(float* out, const float* inp, int N) {
     // Implement this
-    dim3 block(32,1,1);
+    dim3 block(256,1,1);
     dim3 grid((N + block.x -1) / block.x);
 
     gelu_forward_kernel<<<grid, block>>>(out, inp, N);

kernels/gpt2.cuh

@@ -0,0 +1,420 @@
+//-----------------------------------------------------------------------------------------------
+// Copyright (c) 2024 Andrej Karpathy
+// Licensed under the MIT License. See the LICENSE file for details.
+//
+// Modifications Copyright (c) 2025 Hanwen Liu, Hrishi Shah, Kelin Zeng, Charles Pei, and Vijay Daita, ALL RIGHTS RESERVED.
+//-----------------------------------------------------------------------------------------------
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <time.h>
+#include <assert.h>
+#include <float.h>
+#include <string.h>
+#include <unistd.h>
+
+// GPU / CUDA related
+#include <cublas_v2.h>
+#include <cuda_runtime.h>
+#include "utils/utils.h"
+#include "utils/tokenizer.h"
+#include "utils/cuda_utils.cuh"
+
+// kernels 
+#include "kernels/attention.cuh"
+#include "kernels/encoder.cuh"
+#include "kernels/gelu.cuh"
+#include "kernels/layernorm.cuh"
+#include "kernels/matmul.cuh"
+#include "kernels/residual.cuh"
+
+// ----------------------------------------------------------------------------
+// GPT-2 model definition
+
+typedef struct {
+    int max_seq_len;
+    int vocab_size;
+    int padded_vocab_size;  // padded to multiple of 128
+    int num_layers;
+    int num_heads;
+    int channels;
+} GPT2Config;
+
+// the parameters of the model
+#define NUM_PARAMETER_TENSORS 16
+typedef struct {
+    float* wte; // (Vp, C)
+    float* wpe; // (maxT, C)
+    float* ln1w; // (L, C)
+    float* ln1b; // (L, C)
+    float* qkvw; // (L, 3*C, C)
+    float* qkvb; // (L, 3*C)
+    float* attprojw; // (L, C, C)
+    float* attprojb; // (L, C)
+    float* ln2w; // (L, C)
+    float* ln2b; // (L, C)
+    float* fcw; // (L, 4*C, C)
+    float* fcb; // (L, 4*C)
+    float* fcprojw; // (L, C, 4*C)
+    float* fcprojb; // (L, C)
+    float* lnfw; // (C)
+    float* lnfb; // (C)
+} ParameterTensors;
+
+void fill_in_parameter_sizes(size_t* param_sizes, GPT2Config config) {
+    int Vp = config.padded_vocab_size;
+    int C = config.channels;
+    int maxT = config.max_seq_len;
+    int L = config.num_layers;
+    param_sizes[0] = Vp * C; // wte
+    param_sizes[1] = maxT * C; // wpe
+    param_sizes[2] = L * C; // ln1w
+    param_sizes[3] = L * C; // ln1b
+    param_sizes[4] = L * (3 * C) * C; // qkvw
+    param_sizes[5] = L * (3 * C); // qkvb
+    param_sizes[6] = L * C * C; // attprojw
+    param_sizes[7] = L * C; // attprojb
+    param_sizes[8] = L * C; // ln2w
+    param_sizes[9] = L * C; // ln2b
+    param_sizes[10] = L * (4 * C) * C; // fcw
+    param_sizes[11] = L * (4 * C); // fcb
+    param_sizes[12] = L * C * (4 * C); // fcprojw
+    param_sizes[13] = L * C; // fcprojb
+    param_sizes[14] = C; // lnfw
+    param_sizes[15] = C; // lnfb
+}
+
+// allocate memory for the parameters and point the individual tensors to the right places
+float* malloc_and_point_parameters(ParameterTensors* params, size_t* param_sizes, int on_device) {
+    // on_device: 0 = CPU, 1 = GPU
+    // calculate the number of parameters
+    size_t num_parameters = 0;
+    for (size_t i = 0; i < NUM_PARAMETER_TENSORS; i++) {
+        num_parameters += param_sizes[i];
+    }
+    // malloc all parameters all at once on the device
+    float* params_memory;
+    if (on_device) {
+        cudaCheck(cudaMalloc((void**)&params_memory, num_parameters * sizeof(float)));
+    } else {
+        params_memory = (float*)mallocCheck(num_parameters * sizeof(float));
+    }
+    // assign all the tensors their place in the array
+    float** ptrs[] = {
+        &params->wte, &params->wpe, &params->ln1w, &params->ln1b, &params->qkvw, &params->qkvb,
+        &params->attprojw, &params->attprojb, &params->ln2w, &params->ln2b, &params->fcw, &params->fcb,
+        &params->fcprojw, &params->fcprojb, &params->lnfw, &params->lnfb
+    };
+    float* params_memory_iterator = params_memory;
+    for (size_t i = 0; i < NUM_PARAMETER_TENSORS; i++) {
+        *(ptrs[i]) = params_memory_iterator;
+        params_memory_iterator += param_sizes[i];
+    }
+    return params_memory;
+}
+
+#define NUM_ACTIVATION_TENSORS 21
+typedef struct {
+    float* encoded; // (B, T, C)
+    float* ln1; // (L, B, T, C)
+    float* ln1_mean; // (L, B, T)
+    float* ln1_rstd; // (L, B, T)
+    float* atty; // (L, B, T, C)
+    float* att; // (L, B, NH, T, T)
+    float* attproj; // (L, B, T, C)
+    float* residual2; // (L, B, T, C)
+    float* ln2; // (L, B, T, C)
+    float* ln2_mean; // (L, B, T)
+    float* ln2_rstd; // (L, B, T)
+    float* fch; // (L, B, T, 4*C)
+    float* fch_gelu; // (L, B, T, 4*C)
+    float* fcproj; // (L, B, T, C)
+    float* residual3; // (L, B, T, C)
+    float* lnf; // (B, T, C)
+    float* lnf_mean; // (B, T)
+    float* lnf_rstd; // (B, T)
+    float* losses; // (B, T)
+    float* qkvr; // (L, B, T, 3*C)
+    // This buffer will store the logits
+    // during the processing of transformer blocks, we will also use this as a
+    // general scratchpad buffer. Allocation is made large enough to hold (B, T, 3C),
+    // (B, NH, T, T), and (B, T, V) shaped tensors.
+    float* output;
+} ActivationTensors;
+
+void fill_in_activation_sizes(size_t* act_sizes, int B, int T, GPT2Config config) {
+    size_t Vp = config.padded_vocab_size;
+    size_t L = config.num_layers;
+    size_t NH = config.num_heads;
+    size_t C = config.channels;
+    act_sizes[0] = B * T * C; // encoded
+    act_sizes[1] = L * B * T * C; // ln1
+    //to align making 1 to 4
+    act_sizes[2] = 1; // ln1_mean, reserved for future use
+    act_sizes[3] = 1; // ln1_rstd, reserved for future use
+    act_sizes[4] = L * B * T * C; // atty
+    act_sizes[5] = L * B * NH * T * T; // att
+    act_sizes[6] = L * B * T * C; // attproj
+    act_sizes[7] = L * B * T * C; // residual2
+    act_sizes[8] = L * B * T * C; // ln2
+    act_sizes[9] = 1; // ln2_mean, reserved for future use
+    act_sizes[10] = 1; // ln2_rstd, reserved for future use
+    act_sizes[11] = 1; // fch
+    act_sizes[12] = L * B * T * 4*C; // fch_gelu
+    act_sizes[13] = L * B * T * C; // fcproj
+    act_sizes[14] = L * B * T * C; // residual3
+    act_sizes[15] = B * T * C; // lnf
+    act_sizes[16] = 1; // lnf_mean, reserved for future use
+    act_sizes[17] = 1; // lnf_rstd, reserved for future use
+    act_sizes[18] = B * T; // losses
+    act_sizes[19] = L * B * T * 3*C; // qkvr
+    act_sizes[20] = B * T * max(3*C, max(NH*T, Vp)); // output / scratch
+}
+
+// Reserved for future use
+#define NUM_BACKWARD_TENSORS 3
+typedef struct {
+    float* bt4c; // (B, T, 4*C)
+    float* preatt; // (B, NH, T, T)
+    float* residual3; // (B, T, C)
+} GradActTensors;
+
+// op_13: 256-byte aligned activations.
+// Each activation tensor's stride is rounded up to a multiple of 64 floats (256 bytes).
+#define align_to_256_bytes(x) \
+    (x + (128/sizeof(float) - (x % (128/sizeof(float)))) % (128/sizeof(float)))
+
+float* malloc_and_point(float** targets[], const size_t* act_sizes, int n) {
+    size_t num_activations = 0;
+    for (size_t i = 0; i < n; i++) {
+        num_activations += align_to_256_bytes(act_sizes[i]);
+    }
+    float* acts_memory;
+    cudaCheck(cudaMalloc((void**)&acts_memory, num_activations * sizeof(float)));
+    cudaCheck(cudaMemset(acts_memory, 0, num_activations * sizeof(float)));
+    float* acts_memory_iterator = acts_memory;
+    for (size_t i = 0; i < n; i++) {
+        *(targets[i]) = acts_memory_iterator;
+        acts_memory_iterator += align_to_256_bytes(act_sizes[i]);
+    }
+    return acts_memory;
+}
+
+float* malloc_and_point_activations(ActivationTensors* acts, const size_t* act_sizes) {
+    float** ptrs[] = {
+        &acts->encoded, &acts->ln1, &acts->ln1_mean, &acts->ln1_rstd, &acts->atty,
+        &acts->att, &acts->attproj, &acts->residual2, &acts->ln2, &acts->ln2_mean,
+        &acts->ln2_rstd, &acts->fch, &acts->fch_gelu, &acts->fcproj, &acts->residual3, &acts->lnf,
+        &acts->lnf_mean, &acts->lnf_rstd, &acts->losses, &acts->qkvr, &acts->output
+    };
+    return malloc_and_point(ptrs, act_sizes, NUM_ACTIVATION_TENSORS);
+}
+
+typedef struct {
+    GPT2Config config;
+    // the weights of the model, and their sizes
+    ParameterTensors params;
+    size_t param_sizes[NUM_PARAMETER_TENSORS];
+    float* params_memory;
+    size_t num_parameters;
+    ParameterTensors grads;
+    float* grads_memory;
+    float* m_memory;
+    float* v_memory;
+    // the activations of the model, and their sizes
+    ActivationTensors acts;
+    size_t act_sizes[NUM_ACTIVATION_TENSORS];
+    float* acts_memory;
+    size_t num_activations;
+    GradActTensors grads_acts;
+    size_t num_grad_acts;
+    float* grads_acts_memory;
+    // other run state configuration
+    int batch_size; // the batch size (B) of current forward pass
+    int seq_len; // the sequence length (T) of current forward pass
+    int* inputs; // the input tokens for the current forward pass
+    int* targets; // the target tokens for the current forward pass
+    float mean_loss;
+    float* cpu_losses;
+} GPT2;
+
+void gpt2_build_from_checkpoint(GPT2 *model, const char* checkpoint_path) {
+
+    // read in model from a checkpoint file
+    FILE *model_file = fopenCheck(checkpoint_path, "rb");
+    int model_header[256];
+    freadCheck(model_header, sizeof(int), 256, model_file);
+    if (model_header[0] != 20240326) { fprintf(stderr, "GPT2 - Bad magic model file\n"); exit(EXIT_FAILURE); }
+    if (model_header[1] != 3) {
+        // was bumped from 1 -> 3 to incorporate the padded vocab size
+        fprintf(stderr, "GPT2 - Bad version in model file\n");
+        exit(EXIT_FAILURE);
+    }
+
+    // read in hyperparameters
+    model->config.max_seq_len = model_header[2];
+    model->config.vocab_size = model_header[3];
+    model->config.num_layers = model_header[4];
+    model->config.num_heads = model_header[5];
+    model->config.channels = model_header[6];
+    model->config.padded_vocab_size = model_header[7];
+
+    // allocate space for all the parameters and read them in
+    fill_in_parameter_sizes(model->param_sizes, model->config);
+
+    // count the number of parameters
+    size_t num_parameters = 0;
+    for (size_t i = 0; i < NUM_PARAMETER_TENSORS; i++) {
+        num_parameters += model->param_sizes[i];
+    }
+    model->num_parameters = num_parameters;
+
+    // create memory for model parameters on the device
+    model->params_memory = malloc_and_point_parameters(&model->params, model->param_sizes, 1);
+
+    // read in all the parameters from file and copy them to device
+    float* params_memory_cpu = (float*)mallocCheck(num_parameters * sizeof(float));
+    freadCheck(params_memory_cpu, sizeof(float), num_parameters, model_file);
+    cudaCheck(cudaMemcpy(model->params_memory, params_memory_cpu, num_parameters * sizeof(float), cudaMemcpyHostToDevice));
+    free(params_memory_cpu);
+    fcloseCheck(model_file);
+
+    // other inits
+    model->acts_memory = NULL;
+    model->grads_memory = NULL;
+    model->m_memory = NULL;
+    model->v_memory = NULL;
+    model->grads_acts_memory = NULL;
+    model->inputs = NULL;
+    model->targets = NULL;
+    model->cpu_losses = NULL;
+    model->batch_size = 0;
+    model->seq_len = 0;
+    model->mean_loss = -1.0f; // -1.0f will designate no loss
+}
+
+void gpt2_forward(GPT2 *model, int* inputs, int B, int T) {
+
+    // ensure the model was initialized or error out
+    if (model->params_memory == NULL) {
+        printf("GPT2 - Error: model was not initialized properly.\n");
+        exit(EXIT_FAILURE);
+    }
+
+    // convenience parameters
+    int V = model->config.vocab_size;
+    int Vp = model->config.padded_vocab_size;
+    int L = model->config.num_layers;
+    int NH = model->config.num_heads;
+    int C = model->config.channels;
+
+    // validate inputs, all indices must be in the range [0, V)
+    for(int i = 0; i < B * T; i++) {
+        assert(0 <= inputs[i] && inputs[i] < V);
+    }
+
+    // allocate space for all the activations if needed (done here, lazily)
+    if(model->acts_memory == NULL || B != model->batch_size || T != model->seq_len) {
+        // record the current B,T as well
+        model->batch_size = B;
+        model->seq_len = T;
+        // and now allocate the space
+        fill_in_activation_sizes(model->act_sizes, B, T, model->config);
+        size_t num_activations = 0;
+        for (size_t i = 0; i < NUM_ACTIVATION_TENSORS; i++) {
+            num_activations += model->act_sizes[i];
+        }
+        model->num_activations = num_activations;
+        if (model->acts_memory != NULL) {
+            printf("GPT2 - Reallocating activations for B=%d T=%d\n", B, T);
+            cudaFree(model->acts_memory);
+            cudaFree(model->inputs);
+        }
+        model->acts_memory = malloc_and_point_activations(&model->acts, model->act_sizes);
+        printf("GPT2 - Allocated %zu MiB for activations\n", (num_activations * sizeof(float)) >> 20); // >> 20 is /(1024*1024)
+        cudaCheck(cudaMalloc((void**)&model->inputs, B * T * sizeof(int)));
+        cudaCheck(cudaMallocHost((void**)&model->cpu_losses, B * T * sizeof(float)));
+    }
+
+    // copy inputs to the model
+    cudaCheck(cudaMemcpy(model->inputs, inputs, B * T * sizeof(int), cudaMemcpyHostToDevice));
+
+    // forward pass
+    ParameterTensors params = model->params; // for brevity
+    ActivationTensors acts = model->acts;
+    float* residual;
+    encoder_forward(acts.encoded, model->inputs, params.wte, params.wpe, B, T, C); // encoding goes into residual[0]
+
+    for (int l = 0; l < L; l++) {
+
+        residual = l == 0 ? acts.encoded : acts.residual3 + (l-1) * B * T * C;
+
+        // get the pointers of the weights for this layer
+        float* l_ln1w = params.ln1w + l * C;
+        float* l_ln1b = params.ln1b + l * C;
+        float* l_qkvw = params.qkvw + l * 3*C * C;
+        float* l_qkvb = params.qkvb + l * 3*C;
+        float* l_attprojw = params.attprojw + l * C * C;
+        float* l_attprojb = params.attprojb + l * C;
+        float* l_ln2w = params.ln2w + l * C;
+        float* l_ln2b = params.ln2b + l * C;
+        float* l_fcw = params.fcw + l * 4*C * C;
+        float* l_fcb = params.fcb + l * 4*C;
+        float* l_fcprojw = params.fcprojw + l * C * 4*C;
+        float* l_fcprojb = params.fcprojb + l * C;
+
+        // get the pointers of the activations for this layer
+        float* l_ln1 = acts.ln1 + l * B * T * C;
+        float* l_ln1_mean = acts.ln1_mean;
+        float* l_ln1_rstd = acts.ln1_rstd;
+        float* l_qkvr = acts.qkvr + l * B * T * 3*C;
+        float* l_atty = acts.atty + l * B * T * C;
+        float* l_att = acts.att + l * B * NH * T * T;
+        float* l_attproj = acts.attproj + l * B * T * C;
+        float* l_residual2 = acts.residual2 + l * B * T * C;
+        float* l_ln2 = acts.ln2 + l * B * T * C;
+        float* l_ln2_mean = acts.ln2_mean;
+        float* l_ln2_rstd = acts.ln2_rstd;
+        float* l_fch = acts.fch + l * B * T * 4*C;
+        float* l_fch_gelu = acts.fch_gelu + l * B * T * 4*C;
+        float* l_fcproj = acts.fcproj + l * B * T * C;
+        float* l_residual3 = acts.residual3 + l * B * T * C;
+        // these are only needed as scratchpads for the forward pass
+        float* scratch = acts.output;
+
+        // now do the forward pass
+        layernorm_forward(l_ln1, l_ln1_mean, l_ln1_rstd, residual, l_ln1w, l_ln1b, B, T, C);
+        matmul_forward(scratch, l_ln1, l_qkvw, l_qkvb, B, T, C, 3*C);
+        attention_forward(l_atty, l_qkvr, l_att, scratch, B, T, C, NH);
+        matmul_forward(l_attproj, l_atty, l_attprojw, l_attprojb, B, T, C, C);
+        residual_forward(l_residual2, residual, l_attproj, B*T*C);
+        layernorm_forward(l_ln2, l_ln2_mean, l_ln2_rstd, l_residual2, l_ln2w, l_ln2b, B, T, C);
+        //op16 matmul + bias + GeLU fusion
+        // matmul_forward(l_fch, l_ln2, l_fcw, l_fcb, B, T, C, 4*C);
+        // gelu_forward(l_fch_gelu, l_fch, B*T*4*C);
+        matmul_forward_fused_gelu(l_fch_gelu, l_ln2, l_fcw, l_fcb, B, T, C, 4*C);
+        matmul_forward(l_fcproj, l_fch_gelu, l_fcprojw, l_fcprojb, B, T, 4*C, C);
+        residual_forward(l_residual3, l_residual2, l_fcproj, B*T*C);
+    }
+
+    residual = acts.residual3 + (L-1) * B * T * C; // last residual is in residual3
+    layernorm_forward(acts.lnf, acts.lnf_mean, acts.lnf_rstd, residual, params.lnfw, params.lnfb, B, T, C);
+    matmul_forward(acts.output, acts.lnf, params.wte, NULL, B, T, C, Vp);
+
+    model->mean_loss = -1.0f;
+}
+
+void gpt2_free(GPT2 *model) {
+    cudaCheck(cudaFree(model->params_memory));
+    cudaCheck(cudaFree(model->grads_memory));
+    cudaCheck(cudaFree(model->m_memory));
+    cudaCheck(cudaFree(model->v_memory));
+    cudaCheck(cudaFree(model->acts_memory));
+    cudaCheck(cudaFree(model->grads_acts_memory));
+    cudaCheck(cudaFree(model->inputs));
+    cudaCheck(cudaFree(model->targets));
+    cudaFreeHost(model->cpu_losses);
+}
+
+#define GPT2_EOT 50256

kernels/layernorm.cuh

@@ -6,32 +6,88 @@
 #include <float.h>
 #include "../utils/cuda_utils.cuh"
 
+#define BLOCK_DIM 256
+#define WARP_SIZE 32
+
+//lecture 15 slide ~32 
+__device__ __inline__ float warp_reduce(float val) {
+    float partialSum = val;
+    for(unsigned int stride = WARP_SIZE/2; stride > 0; stride /= 2) {
+        partialSum += __shfl_down_sync(0xffffffff, partialSum, stride);
+    }
+    return partialSum;
+}
+
 __global__ void layernorm_forward_kernel(float* out, float* mean, float* rstd, const float* inp, const float* weight,
                                          const float* bias, int B, int T, int C) {
     // Implement this
-    int b = blockDim.x * blockIdx.x + threadIdx.x;
+    int bt = blockIdx.x;
-    int t = blockDim.y * blockIdx.y + threadIdx.y;
+    if(bt >= B*T){
-
-    if (b >= B || t >= T) {
         return;
     }
 
-    float sum = 0.0f;
+    const float* x = inp + bt * C;
-    for (int c = 0; c < C; ++c) {
+    float* o = out + bt * C;
-        sum += inp[b * T * C + t * C + c];
-    }
-    float m = sum / C;
 
-    sum = 0.0f;
+    // computer mean
-    for (int c = 0; c < C; ++c) {
+    float partial_sum = 0.0f;
-        float diff = inp[b * T * C + t * C + c] - m;
+    for(int c = threadIdx.x ; c < C ; c += BLOCK_DIM){
-        sum += diff * diff;
+        partial_sum += x[c];
     }
-    float s = rsqrtf(sum / C + 1e-5f);
+
-    
+    partial_sum = warp_reduce(partial_sum);
-    for (int c = 0; c < C; ++c) {
+
-        float n = s * (inp[b * T * C + t * C + c] - m);
+    __shared__ float partial_sum_s[BLOCK_DIM / WARP_SIZE];
-        out[b * T * C + t * C + c] = n * weight[c] + bias[c];
+    if(threadIdx.x % WARP_SIZE == 0){
+        partial_sum_s[threadIdx.x / WARP_SIZE] = partial_sum;
+    }
+
+    __syncthreads();
+
+    if(threadIdx.x < WARP_SIZE){
+        partial_sum = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_sum_s[threadIdx.x] : 0.0f;
+        partial_sum = warp_reduce(partial_sum);
+    }
+
+    //thread 0 will have the final sum
+    __shared__ float s_mean;
+    if (threadIdx.x == 0) {
+        s_mean = partial_sum / C;
+    }
+
+    __syncthreads();
+    float m = s_mean;
+
+    //compute variance
+    partial_sum = 0.0f;
+    for (int c = threadIdx.x; c < C; c += BLOCK_DIM) {
+        float diff = x[c] - m;
+        partial_sum += diff * diff;
+    }
+
+    partial_sum = warp_reduce(partial_sum);
+
+    if (threadIdx.x % WARP_SIZE == 0) {
+        partial_sum_s[threadIdx.x / WARP_SIZE] = partial_sum;
+    }
+    __syncthreads();
+
+    if (threadIdx.x < WARP_SIZE) {
+        partial_sum = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_sum_s[threadIdx.x] : 0.0f;
+        partial_sum = warp_reduce(partial_sum);
+    }
+
+    __shared__ float s_rstd;
+    if (threadIdx.x == 0) {
+        s_rstd = rsqrtf(partial_sum / C + 1e-5f);
+    }
+    __syncthreads();
+    float s = s_rstd;
+
+    //normalize
+    for (int c = threadIdx.x; c < C; c += BLOCK_DIM) {
+        float n = s * (x[c] - m);
+        o[c] = n * weight[c] + bias[c];
     }
 }
 
@@ -39,10 +95,8 @@ __global__ void layernorm_forward_kernel(float* out, float* mean, float* rstd, c
 void layernorm_forward(float* out, float* mean, float* rstd, float* inp, float* weight, float* bias,
                        int B, int T, int C) {
     // Implement this
-    dim3 block(16,16);
+    int num_blocks = B*T;
-    dim3 grid((B+block.x-1)/block.x, (T+block.y-1)/block.y);
+    layernorm_forward_kernel<<<num_blocks, BLOCK_DIM>>>(out,mean, rstd, inp, weight, bias,B,T, C);
-
-    layernorm_forward_kernel<<<grid, block>>>(out,mean, rstd, inp, weight, bias,B,T, C);
 }
 
 #endif // __LAYERNORM_KERNEL_CUH__

kernels/matmul.cuh

@@ -3,35 +3,213 @@
 
 #include <cuda_runtime.h>
 #include "../utils/cuda_utils.cuh"
+#include "cutlass/cutlass.h"
+#include "cutlass/gemm/device/gemm.h"
+#include "/work/hdd/bche/hshah6/cutlass/include/cutlass/epilogue/thread/linear_combination_gelu.h"
 
-__global__ void matmul_forward_kernel(float* out, const float* inp, const float* weight, 
+//referenced: https://github.com/NVIDIA/cutlass/blob/main/examples/14_ampere_tf32_tensorop_gemm/ampere_tf32_tensorop_gemm.cu
-                                      const float* bias,int B,int T,  int C, int OC) {
-    // Implement this
-    int b = blockDim.x*blockIdx.x + threadIdx.x; // B
-    int t = blockDim.y*blockIdx.y + threadIdx.y; // T
-    int oc = blockDim.z*blockIdx.z + threadIdx.z; // OC
 
-    if (b >= B || t >= T || oc >= OC) {
+//removing bias overhead and fusing with cutlass
-        return;
+// __global__ void bias_add_filler_kernel(float* __restrict__ out, const float* __restrict__ bias, int BT, int OC) {
+//     int x = blockIdx.x * blockDim.x + threadIdx.x;
+//     if (x < BT * OC) out[x] = bias[x % OC];
+// }
+
+//using big cutlass for kernels for BT>=64
+using GemmDefault = cutlass::gemm::device::Gemm<
+    float, cutlass::layout::RowMajor,
+    float, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor,
+    float,
+    cutlass::arch::OpClassTensorOp,
+    cutlass::arch::Sm80,
+    cutlass::gemm::GemmShape<128,128,32>,
+    cutlass::gemm::GemmShape<64,64,32>,
+    cutlass::gemm::GemmShape<16,8,8>,
+    cutlass::epilogue::thread::LinearCombination<float, 128 / cutlass::sizeof_bits<float>::value, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
+    3 // Stages
+>;
+
+//small cutlass for kernels for BT<64, we can use smaller threadblocks and warps to increase occupancy
+using GemmsmallDefault = cutlass::gemm::device::Gemm<
+    float, cutlass::layout::RowMajor,
+    float, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor,
+    float,
+    cutlass::arch::OpClassTensorOp,
+    cutlass::arch::Sm80,
+    cutlass::gemm::GemmShape<32,128,32>,
+    cutlass::gemm::GemmShape<32,64,32>,
+    cutlass::gemm::GemmShape<16,8,8>,
+    cutlass::epilogue::thread::LinearCombination<float, 128 / cutlass::sizeof_bits<float>::value, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
+    2
+>;
+
+using GemmMediumDefault = cutlass::gemm::device::Gemm<
+    float, cutlass::layout::RowMajor,
+    float, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor,
+    float,
+    cutlass::arch::OpClassTensorOp,
+    cutlass::arch::Sm80,
+    cutlass::gemm::GemmShape<64,128,32>,
+    cutlass::gemm::GemmShape<32,64,32>,
+    cutlass::gemm::GemmShape<16,8,8>,
+    cutlass::epilogue::thread::LinearCombination<float, 128 / cutlass::sizeof_bits<float>::value, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
+    3
+>;
+
+inline void matmul_forward(float* __restrict__ out, const float* __restrict__ inp, const float* __restrict__ weight, 
+                                      const float* __restrict__ bias,int B,int T,  int C, int OC) {
+    int BT = B*T;
+
+    if(bias != NULL){
+        if(BT <= 32){
+            GemmsmallDefault gemm_op;
+            typename GemmsmallDefault::Arguments args{ 
+                {BT, OC, C}, // problem size
+                {inp, C}, // A matrix
+                {weight, C}, // B matrix
+                {bias, 0}, // C matrix
+                {out, OC}, // D matrix (output)
+                {1.0f, 1.0f} // alpha, beta for the epilogue
+            };
+            gemm_op(args);
+        }
+        else if(BT <= 128){
+            GemmMediumDefault gemm_op;
+            typename GemmMediumDefault::Arguments args{ 
+                {BT, OC, C}, // problem size
+                {inp, C}, // A matrix
+                {weight, C}, // B matrix
+                {bias, 0}, // C matrix
+                {out, OC}, // D matrix (output)
+                {1.0f, 1.0f} // alpha, beta for the epilogue
+            };
+            gemm_op(args);
+        }
+        else {
+            GemmDefault gemm_op;
+            typename GemmDefault::Arguments args{ 
+                {BT, OC, C}, // problem size
+                {inp, C}, // A matrix
+                {weight, C}, // B matrix
+                {bias, 0}, // C matrix
+                {out, OC}, // D matrix (output)
+                {1.0f, 1.0f} // alpha, beta for the epilogue
+            };
+            gemm_op(args);
+        }
+    }
+    else{
+        if(BT <= 32){
+            GemmsmallDefault gemm_op;
+            typename GemmsmallDefault::Arguments args{ 
+                {BT, OC, C}, // problem size
+                {inp, C}, // A matrix
+                {weight, C}, // B matrix
+                {out, OC}, // C matrix (output)
+                {out, OC}, // D matrix (output)
+                {1.0f, 0.0f} // alpha, beta for the epilogue
+            };
+            gemm_op(args);
+        }
+        else if(BT <= 128){
+            GemmDefault gemm_op;
+            typename GemmDefault::Arguments args{ 
+                {BT, OC, C}, // problem size
+                {inp, C}, // A matrix
+                {weight, C}, // B matrix
+                {out, OC}, // C matrix (output)
+                {out, OC}, // D matrix (output)
+                {1.0f, 0.0f} // alpha, beta for the epilogue
+            };
+            gemm_op(args);
+        }
+        else{
+            GemmDefault gemm_op;
+            typename GemmDefault::Arguments args{ 
+                {BT, OC, C}, // problem size
+                {inp, C}, // A matrix
+                {weight, C}, // B matrix
+                {out, OC}, // C matrix (output)
+                {out, OC}, // D matrix (output)
+                {1.0f, 0.0f} // alpha, beta for the epilogue
+            };
+            gemm_op(args);
+        }
     }
 
-    float sum = (bias != NULL) ? bias[oc] : 0.0f;
+    
-    const float* inp_token_ptr = inp + b * T * C + t * C;
-    for(int i = 0; i < C; i++){
-        sum += inp_token_ptr[i] * weight[oc*C + i] ;
-    }
-    out[b * T * OC + t * OC + oc] = sum;
 }
 
+using Gemmgelu = cutlass::gemm::device::Gemm<
+    float, cutlass::layout::RowMajor,
+    float, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor,
+    float,
+    cutlass::arch::OpClassTensorOp,
+    cutlass::arch::Sm80,
+    cutlass::gemm::GemmShape<128, 128, 32>,
+    cutlass::gemm::GemmShape<64, 64, 32>,
+    cutlass::gemm::GemmShape<16, 8, 8>,
+    cutlass::epilogue::thread::LinearCombinationGELU<float, 128 / cutlass::sizeof_bits<float>::value, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
+    3>;
+
+using GemmgeluSmall = cutlass::gemm::device::Gemm<
+    float, cutlass::layout::RowMajor,
+    float, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor,
+    float,
+    cutlass::arch::OpClassTensorOp,
+    cutlass::arch::Sm80,
+    cutlass::gemm::GemmShape<32, 128, 32>,
+    cutlass::gemm::GemmShape<32, 64, 32>,
+    cutlass::gemm::GemmShape<16, 8, 8>,
+    cutlass::epilogue::thread::LinearCombinationGELU<float, 128 / cutlass::sizeof_bits<float>::value, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
+    2>;
+
+using GemmgeluMedium = cutlass::gemm::device::Gemm<
+    float, cutlass::layout::RowMajor,
+    float, cutlass::layout::ColumnMajor,
+    float, cutlass::layout::RowMajor,
+    float,
+    cutlass::arch::OpClassTensorOp,
+    cutlass::arch::Sm80,
+    cutlass::gemm::GemmShape<64, 128, 32>,
+    cutlass::gemm::GemmShape<32, 64, 32>,
+    cutlass::gemm::GemmShape<16, 8, 8>,
+    cutlass::epilogue::thread::LinearCombinationGELU<float, 128 / cutlass::sizeof_bits<float>::value, float, float>,
+    cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>,
+    3>;
+
 // Launch kernel here
-void matmul_forward(float* out, const float* inp, const float* weight, const float* bias,
+inline void matmul_forward_fused_gelu(float* out, const float* inp, const float* weight, const float* bias,
                     int B, int T, int C, int OC) {
     // Implement this
-    dim3 block(8,8,8);
+    int BT = B*T;
-    dim3 grid((B+block.x-1)/block.x, (T+block.y-1)/block.y, (OC+block.z-1)/block.z);
-
-    matmul_forward_kernel<<<grid, block>>>(out, inp, weight, bias,B,T, C, OC);
 
+    if (BT <= 32) {
+        GemmgeluSmall op;
+        typename GemmgeluSmall::Arguments args{
+            {BT, OC, C}, {inp, C}, {weight, C}, {bias, 0}, {out, OC}, {1.0f, 1.0f}};
+        op(args);
+    } else if (BT <= 128) {
+        GemmgeluMedium op;
+        typename GemmgeluMedium::Arguments args{
+            {BT, OC, C}, {inp, C}, {weight, C}, {bias, 0}, {out, OC}, {1.0f, 1.0f}};
+        op(args);
+    } else {
+        Gemmgelu op;
+        typename Gemmgelu::Arguments args{
+            {BT, OC, C}, {inp, C}, {weight, C}, {bias, 0}, {out, OC}, {1.0f, 1.0f}};
+        op(args);
+    }
 }
 
 #endif // __MATMUL_KERNEL_CUH__

kernels/residual.cuh

@@ -4,7 +4,7 @@
 #include <cuda_runtime.h>
 #include "../utils/cuda_utils.cuh"
 
-__global__ void residual_forward_kernel(float* out, float* inp1, float* inp2, int N) {
+__global__ void residual_forward_kernel(float* __restrict__ out, const float* __restrict__ inp1, const float* __restrict__ inp2, int N) {
     // Implement this
     int x = blockDim.x * blockIdx.x + threadIdx.x;
 
@@ -16,7 +16,7 @@ __global__ void residual_forward_kernel(float* out, float* inp1, float* inp2, in
 // Launch kernel here
 void residual_forward(float* out, float* inp1, float* inp2, int N) {
     // Implement this
-    dim3 block(32,1,1);
+    dim3 block(256,1,1);
     dim3 grid((N + block.x - 1)/ block.x);
 
     residual_forward_kernel<<<grid, block>>>(out, inp1, inp2, N);

kernels/softmax.cuh

@@ -6,30 +6,93 @@
 #include <float.h>
 #include "../utils/cuda_utils.cuh"
 
+#define BLOCK_DIM 256
+#define WARP_SIZE 32
+
+// lecture 15 slide 32
+__device__ __inline__ float warpReduceSum(float partialSum) {
+    for (unsigned int stride = WARP_SIZE / 2; stride > 0; stride /= 2) {
+        partialSum += __shfl_down_sync(0xffffffff, partialSum, stride);
+    }
+    return partialSum;
+}
+
+// same pattern for max
+__device__ __inline__ float warpReduceMax(float partialMax) {
+    for (unsigned int stride = WARP_SIZE / 2; stride > 0; stride /= 2) {
+        partialMax = fmaxf(partialMax, __shfl_down_sync(0xffffffff, partialMax, stride));
+    }
+    return partialMax;
+}
+
 __global__ void softmax_forward_kernel(float* out, float inv_temperature, const float* inp, int N, int T) {
     // Implement this
-    int x = blockDim.x * blockIdx.x + threadIdx.x;
+    int row = blockIdx.x;
+    if(row >= N*T){
+        return;
+    }
 
-    if(x < N*T){
+    int own_pos = row % T;
-        int own_pos = x % T;
+    const float* a = inp + row * T;
-        const float* a = inp + x * T;
+    float* o = out + row * T;
 
-        float maxval = -FLT_MAX;
+    //find max
-        for (int i = 0; i <= own_pos; i++) {
+    //thread coarsening
-            maxval = fmaxf(maxval,a[i]);
+    float partial_max = -FLT_MAX;
-        }
+    for(int i = threadIdx.x; i <= own_pos; i+=BLOCK_DIM){
+        partial_max = fmaxf(partial_max,a[i]);
+    }
 
-        float sumval = 0.0f;
+    //lecture 15 slide 38
-        for (int i = 0; i <= own_pos; ++i) {
+    partial_max = warpReduceMax(partial_max);
-            float ev = expf(inv_temperature * (a[i] - maxval));
-            sumval += ev;
-            out[x * T + i] = ev;
-        }
 
-        float norm = 1.0f / sumval;
+    __shared__ float partial_max_s[BLOCK_DIM / WARP_SIZE];
-        for (int i = 0; i <= own_pos; ++i) {
+    if (threadIdx.x % WARP_SIZE == 0) {
-            out[x * T + i] *= norm;
+        partial_max_s[threadIdx.x / WARP_SIZE] = partial_max;
-        }
+    }
+    __syncthreads();
+    
+    if (threadIdx.x < WARP_SIZE) {
+        partial_max = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_max_s[threadIdx.x] : -FLT_MAX;
+        partial_max = warpReduceMax(partial_max);
+    }
+
+    __shared__ float s_max;
+    if (threadIdx.x == 0) {
+        s_max = partial_max;
+    }
+    __syncthreads();
+    float maxval = s_max;
+
+    float partial_sum = 0.0f;
+    for (int i = threadIdx.x; i <= own_pos; i += BLOCK_DIM) {
+        float ev = expf(inv_temperature * (a[i] - maxval));
+        o[i] = ev;
+        partial_sum += ev;
+    }
+
+    partial_sum = warpReduceSum(partial_sum);
+
+    __shared__ float partial_sums_s[BLOCK_DIM / WARP_SIZE];
+    if (threadIdx.x % WARP_SIZE == 0) {
+        partial_sums_s[threadIdx.x / WARP_SIZE] = partial_sum;
+    }
+    __syncthreads();
+
+    if (threadIdx.x < WARP_SIZE) {
+        partial_sum = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_sums_s[threadIdx.x] : 0.0f;
+        partial_sum = warpReduceSum(partial_sum);
+    }
+
+    __shared__ float s_sum;
+    if (threadIdx.x == 0) {
+        s_sum = partial_sum;
+    }
+    __syncthreads();
+    float norm = 1.0f / s_sum;
+
+    for (int i = threadIdx.x; i <= own_pos; i += BLOCK_DIM) {
+        o[i] *= norm;
     }
 }
 

kernels_req_1/attention.cuh

@@ -1,191 +0,0 @@
-#ifndef __ATTENTION_CUH__
-#define __ATTENTION_CUH__
-
-#include <cuda_runtime.h>
-#include "../utils/cuda_utils.cuh"
-#include "softmax.cuh"
-#include <mma.h>
-
-using namespace nvcuda;
-
-#define WMMA_M 16
-#define WMMA_N 16
-#define WMMA_K 8
-
-__global__ void permute_kernel(float* q, float* k, float* v, const float* inp, int B, int N, int NH, int d) {
-    // Implement this
-
-    int x = blockDim.x * blockIdx.x + threadIdx.x;
-
-    if( x < B*NH*N*d){
-        int b = (x / d / N / NH);
-        int nh = (x / d / N) % NH;
-        int t = (x / d) % N;
-        int d_ = x % d;
-
-        int inp_idx = (b * N * 3 * NH * d) + (t * 3 * NH * d) + (0 * NH * d) + (nh * d) + d_;
-        q[b * NH * N * d + nh * N * d + t * d + d_] = inp[inp_idx];
-        k[b * NH * N * d + nh * N * d + t * d + d_] = inp[inp_idx + NH * d];
-        v[b * NH * N * d + nh * N * d + t * d + d_] = inp[inp_idx + 2 * (NH * d)];
-    }
-}
-
-
-__global__ void unpermute_kernel(float* inp, float *out, int B, int N, int NH, int d) {
-    // Implement this
-    int x = blockDim.x * blockIdx.x + threadIdx.x;
-
-    if( x < B*NH*N*d){
-        int b = (x / d / N / NH);
-        int nh = (x / d / N) % NH;
-        int t = (x / d) % N;
-        int d_ = x % d;
-
-        out[b*N*NH*d + t*NH*d + nh*d + d_] = inp[b*NH*N*d + nh*N*d + t*d + d_];
-    }
-
-}
-
-#define PREATT_WARPS_M 2
-#define PREATT_WARPS_N 4
-#define PREATT_BM (PREATT_WARPS_M * WMMA_M)
-#define PREATT_BN (PREATT_WARPS_N * WMMA_N)
-
-__global__ void preattention_scores_kernel(float* q, float* k, int B, int T, int NH, int HS, float* preatt){
-    // each block handles a (b,nh) head and a tile of (t1, t2)
-    int head = blockIdx.z; //whcih pair
-    int b = head / NH;
-    int nh = head % NH;
-
-    //head pointers 
-    const float* q_head = q + (b * NH + nh) * T * HS; //[T,HS]
-    const float* k_head = k + (b * NH + nh) * T * HS; //[T,HS]
-    float* preatt_head = preatt + (b * NH +nh) * T * T; //[T,T]
-
-    int start_row = blockIdx.x * PREATT_BM;
-    int start_col = blockIdx.y * PREATT_BN;
-
-    int warp_id = threadIdx.x / 32;
-    int warp_row = warp_id / PREATT_WARPS_N;
-    int warp_col = warp_id % PREATT_WARPS_N;
-
-    int warp_out_row = start_row + warp_row * WMMA_M;
-    int warp_out_col = start_col + warp_col * WMMA_N;
-
-    // Q * K^T: A=Q (row_major), B=K (col_major gives K^T)
-    wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, wmma::precision::tf32, wmma::row_major> a_frag;
-    wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, wmma::precision::tf32, wmma::col_major> b_frag;
-    wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
-
-    wmma::fill_fragment(c_frag, 0.0f);
-
-    for(int kk = 0; kk < HS; kk += WMMA_K){
-        if(warp_out_row < T && kk < HS){
-            wmma::load_matrix_sync(a_frag, q_head + warp_out_row * HS + kk, HS);
-            #pragma unroll
-            for(int i = 0; i < a_frag.num_elements; i++){
-                a_frag.x[i] = wmma::__float_to_tf32(a_frag.x[i]);
-            }
-        }
-
-        if(warp_out_col < T && kk < HS) {
-            wmma::load_matrix_sync(b_frag, k_head + warp_out_col * HS + kk, HS);
-            #pragma unroll
-            for (int i = 0; i < b_frag.num_elements; i++)
-                b_frag.x[i] = wmma::__float_to_tf32(b_frag.x[i]);
-        }
-
-        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
-    }
-
-    //write to global mem
-    if(warp_out_row < T && warp_out_col < T){
-        wmma::store_matrix_sync(preatt_head + warp_out_row * T + warp_out_col, c_frag, T, wmma::mem_row_major);
-    }
-}
-
-#define ATTV_WARPS_M 2
-#define ATTV_WARPS_N 4
-#define ATTV_BM 32
-#define ATTV_BN 64
-
-__global__ void attention_v_kernel(float* vaccum, float* att, float* v, int B, int T, int NH, int HS){
-    int head = blockIdx.z; //which pair
-    int b = head / NH;
-    int nh = head % NH;
-
-    //head pointers 
-    const float* att_head = att + (b * NH + nh) * T * T; //[T,HS]
-    const float* v_head = v + (b * NH + nh) * T * HS; //[T,HS]
-    float* out_head = vaccum + (b * NH +nh) * T * HS; //[T,T]
-
-    int row_start = blockIdx.x * ATTV_BM;
-    int col_start = blockIdx.y * ATTV_BN;
-
-    int warpId = threadIdx.x / 32;
-    int warp_row = warpId / ATTV_WARPS_N;
-    int warp_col = warpId % ATTV_WARPS_N;
-
-    int warp_out_row = row_start + warp_row * WMMA_M;
-    int warp_out_col = col_start + warp_col * WMMA_N;
-
-    // att * V: both row-major; standard matmul
-    wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, wmma::precision::tf32, wmma::row_major> a_frag;
-    wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, wmma::precision::tf32, wmma::row_major> b_frag;
-    wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
-
-    wmma::fill_fragment(c_frag, 0.0f);
-
-    for(int kk = 0; kk < T; kk+=WMMA_K){
-        if (warp_out_row < T && kk < T) {
-            wmma::load_matrix_sync(a_frag, att_head + warp_out_row * T + kk, T);
-            #pragma unroll
-            for (int i = 0; i < a_frag.num_elements; i++)
-                a_frag.x[i] = wmma::__float_to_tf32(a_frag.x[i]);
-        }
-
-        if (kk < T && warp_out_col < HS) {
-            wmma::load_matrix_sync(b_frag, v_head + kk * HS + warp_out_col, HS);
-            #pragma unroll
-            for (int i = 0; i < b_frag.num_elements; i++)
-                b_frag.x[i] = wmma::__float_to_tf32(b_frag.x[i]);
-        }
-
-        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
-    }
-
-    if (warp_out_row < T && warp_out_col < HS) {
-        wmma::store_matrix_sync(out_head + warp_out_row * HS + warp_out_col, c_frag, HS, wmma::mem_row_major);
-    }
-}
-
-// Launch all kernels related to attention here 
-void attention_forward(float* out, float* qkvr, float* att, float* inp, int B, int T, int C, int NH) {
-    // Implement this
-    int HS = C / NH;
-
-    float *q, *k, *v;
-    q = qkvr + 0 * B * T * C;
-    k = qkvr + 1 * B * T * C;
-    v = qkvr + 2 * B * T * C;
-
-    dim3 block(32,1,1);
-    dim3 grid((B * NH * T * HS + block.x-1) / block.x,1,1);
-    permute_kernel<<<grid, block>>>(q, k, v, inp, B, T, NH, HS);
-
-    float* preatt = inp;
-    dim3 grid_preatt((T + PREATT_BM - 1) / PREATT_BM, (T + PREATT_BN - 1) / PREATT_BN, B * NH);
-    preattention_scores_kernel<<<grid_preatt, 256>>>(q, k, B, T, NH, HS, preatt);
-
-    float scale = 1.0 / sqrtf(HS);
-    dim3 grid_softmax (( B * NH * T + block.x -1 )/ block.x,1,1);
-    softmax_forward_kernel<<<grid_softmax, block>>>(att, scale, preatt, B * NH, T);
-
-    float* vaccum = inp;
-    dim3 grid_attv((T + ATTV_BM - 1) / ATTV_BM, (HS + ATTV_BN - 1) / ATTV_BN, B * NH);
-    attention_v_kernel<<<grid_attv,256>>>(vaccum, att, v, B, T, NH, HS);
-
-    unpermute_kernel<<<grid, block>>>(vaccum, out, B, T, NH, HS);
-}
-
-#endif // __ATTENTION_CUH__

kernels_req_1/matmul.cuh

@@ -1,98 +0,0 @@
-#ifndef __MATMUL_KERNEL_CUH__
-#define __MATMUL_KERNEL_CUH__
-
-#include <cuda_runtime.h>
-#include <mma.h>
-#include "../utils/cuda_utils.cuh"
-
-using namespace nvcuda;
-
-#define WMMA_M 16
-#define WMMA_N 16
-#define WMMA_K 8
-
-#define WARPS_M 4
-#define WARPS_N 2
-#define BM 64
-#define BN 32
-#define BK WMMA_K
-
-__global__ void matmul_forward_kernel(float* out, const float* inp, const float* weight, 
-                                      const float* bias,int B,int T,  int C, int OC) {
-    // Implement this
-    int BT = B*T;
-
-    // block output
-    int start_row = blockIdx.x * BM;
-    int start_col = blockIdx.y * BN;
-
-    // which warp
-    int warp_id = threadIdx.x / 32;
-    int warp_row = warp_id / WARPS_N;
-    int warp_col = warp_id % WARPS_N;
-
-    int warp_out_row = start_row + warp_row * WMMA_M;
-    int warp_out_col = start_col + warp_col * WMMA_N;
-
-    wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, wmma::precision::tf32, wmma::row_major> a_frag;
-    wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, wmma::precision::tf32, wmma::col_major> b_frag;
-    wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
-
-    wmma::fill_fragment(c_frag, 0.0f);
-
-    for(int k = 0; k < C; k += WMMA_K){
-        if(warp_out_row < BT && k < C){
-            wmma::load_matrix_sync(a_frag, inp + warp_out_row * C + k, C);
-
-            #pragma unroll
-            for(int i = 0; i< a_frag.num_elements; i++){
-                a_frag.x[i] = wmma::__float_to_tf32(a_frag.x[i]);
-            }
-        }
-
-        if (warp_out_col < OC && k < C) {
-            wmma::load_matrix_sync(b_frag, weight + warp_out_col * C + k, C);
-            #pragma unroll
-            for (int i = 0; i < b_frag.num_elements; i++) {
-                b_frag.x[i] = wmma::__float_to_tf32(b_frag.x[i]);
-            }
-        }
-
-        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
-    }
-
-    //store  in s mem, add bais durinthe writw out  for performance
-    __shared__ float out_tile[BM * BN];
-
-    int row_smem = warp_row * WMMA_M;
-    int col_smem = warp_col * WMMA_N;
-    wmma::store_matrix_sync(out_tile + row_smem * BN + col_smem, c_frag, BN, wmma::mem_row_major);
-    __syncthreads();
-
-    for(int i = threadIdx.x; i < BM*BN; i += blockDim.x){
-        int r_l = i / BN;
-        int c_l = i % BN;
-        int r_g = start_row + r_l;
-        int c_g = start_col + c_l;
-
-        if(r_g < BT && c_g < OC){
-            if(bias != NULL){
-                out_tile[i] += bias[c_g];
-            }
-            out[r_g * OC + c_g] = out_tile[i];
-        }
-    }
-}
-
-// Launch kernel here
-void matmul_forward(float* out, const float* inp, const float* weight, const float* bias,
-                    int B, int T, int C, int OC) {
-    // Implement this
-    int BT = B*T;
-    dim3 block(256);
-    dim3 grid((BT + BM - 1) / BM, (OC + BN - 1) / BN);
-    matmul_forward_kernel<<<grid, block>>>(out, inp, weight, bias,B,T, C, OC);
-
-}
-
-#endif // __MATMUL_KERNEL_CUH__

kernels_req_3/layernorm.cuh

@@ -1,102 +0,0 @@
-#ifndef __LAYERNORM_KERNEL_CUH__
-#define __LAYERNORM_KERNEL_CUH__
-
-#include <cuda_runtime.h>
-#include <math.h>
-#include <float.h>
-#include "../utils/cuda_utils.cuh"
-
-#define BLOCK_DIM 256
-#define WARP_SIZE 32
-
-//lecture 15 slide ~32 
-__device__ __inline__ float warp_reduce(float val) {
-    float partialSum = val;
-    for(unsigned int stride = WARP_SIZE/2; stride > 0; stride /= 2) {
-        partialSum += __shfl_down_sync(0xffffffff, partialSum, stride);
-    }
-    return partialSum;
-}
-
-__global__ void layernorm_forward_kernel(float* out, float* mean, float* rstd, const float* inp, const float* weight,
-                                         const float* bias, int B, int T, int C) {
-    // Implement this
-    int bt = blockIdx.x;
-    if(bt >= B*T){
-        return;
-    }
-
-    const float* x = inp + bt * C;
-    float* o = out + bt * C;
-
-    // computer mean
-    float partial_sum = 0.0f;
-    for(int c = threadIdx.x ; c < C ; c += BLOCK_DIM){
-        partial_sum += x[c];
-    }
-
-    partial_sum = warp_reduce(partial_sum);
-
-    __shared__ float partial_sum_s[BLOCK_DIM / WARP_SIZE];
-    if(threadIdx.x % WARP_SIZE == 0){
-        partial_sum_s[threadIdx.x / WARP_SIZE] = partial_sum;
-    }
-
-    __syncthreads();
-
-    if(threadIdx.x < WARP_SIZE){
-        partial_sum = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_sum_s[threadIdx.x] : 0.0f;
-        partial_sum = warp_reduce(partial_sum);
-    }
-
-    //thread 0 will have the final sum
-    __shared__ float s_mean;
-    if (threadIdx.x == 0) {
-        s_mean = partial_sum / C;
-    }
-
-    __syncthreads();
-    float m = s_mean;
-
-    //compute variance
-    partial_sum = 0.0f;
-    for (int c = threadIdx.x; c < C; c += BLOCK_DIM) {
-        float diff = x[c] - m;
-        partial_sum += diff * diff;
-    }
-
-    partial_sum = warp_reduce(partial_sum);
-
-    if (threadIdx.x % WARP_SIZE == 0) {
-        partial_sum_s[threadIdx.x / WARP_SIZE] = partial_sum;
-    }
-    __syncthreads();
-
-    if (threadIdx.x < WARP_SIZE) {
-        partial_sum = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_sum_s[threadIdx.x] : 0.0f;
-        partial_sum = warp_reduce(partial_sum);
-    }
-
-    __shared__ float s_rstd;
-    if (threadIdx.x == 0) {
-        s_rstd = rsqrtf(partial_sum / C + 1e-5f);
-    }
-    __syncthreads();
-    float s = s_rstd;
-
-    //normalize
-    for (int c = threadIdx.x; c < C; c += BLOCK_DIM) {
-        float n = s * (x[c] - m);
-        o[c] = n * weight[c] + bias[c];
-    }
-}
-
-// Launch kernel here
-void layernorm_forward(float* out, float* mean, float* rstd, float* inp, float* weight, float* bias,
-                       int B, int T, int C) {
-    // Implement this
-    int num_blocks = B*T;
-    layernorm_forward_kernel<<<num_blocks, BLOCK_DIM>>>(out,mean, rstd, inp, weight, bias,B,T, C);
-}
-
-#endif // __LAYERNORM_KERNEL_CUH__

kernels_req_3/softmax.cuh

@@ -1,99 +0,0 @@
-#ifndef __SOFTMAX_KERNEL_CUH__
-#define __SOFTMAX_KERNEL_CUH__
-
-#include <cuda_runtime.h>
-#include <math.h>
-#include <float.h>
-#include "../utils/cuda_utils.cuh"
-
-#define BLOCK_DIM 256
-#define WARP_SIZE 32
-
-// lecture 15 slide 32
-__device__ __inline__ float warpReduceSum(float partialSum) {
-    for (unsigned int stride = WARP_SIZE / 2; stride > 0; stride /= 2) {
-        partialSum += __shfl_down_sync(0xffffffff, partialSum, stride);
-    }
-    return partialSum;
-}
-
-// same pattern for max
-__device__ __inline__ float warpReduceMax(float partialMax) {
-    for (unsigned int stride = WARP_SIZE / 2; stride > 0; stride /= 2) {
-        partialMax = fmaxf(partialMax, __shfl_down_sync(0xffffffff, partialMax, stride));
-    }
-    return partialMax;
-}
-
-__global__ void softmax_forward_kernel(float* out, float inv_temperature, const float* inp, int N, int T) {
-    // Implement this
-    int row = blockIdx.x;
-    if(row >= N*T){
-        return;
-    }
-
-    int own_pos = row % T;
-    const float* a = inp + row * T;
-    float* o = out + row * T;
-
-    //find max
-    //thread coarsening
-    float partial_max = -FLT_MAX;
-    for(int i = threadIdx.x; i <= own_pos; i+=BLOCK_DIM){
-        partial_max = fmaxf(partial_max,a[i]);
-    }
-
-    //lecture 15 slide 38
-    partial_max = warpReduceMax(partial_max);
-
-    __shared__ float partial_max_s[BLOCK_DIM / WARP_SIZE];
-    if (threadIdx.x % WARP_SIZE == 0) {
-        partial_max_s[threadIdx.x / WARP_SIZE] = partial_max;
-    }
-    __syncthreads();
-    
-    if (threadIdx.x < WARP_SIZE) {
-        partial_max = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_max_s[threadIdx.x] : -FLT_MAX;
-        partial_max = warpReduceMax(partial_max);
-    }
-
-    __shared__ float s_max;
-    if (threadIdx.x == 0) {
-        s_max = partial_max;
-    }
-    __syncthreads();
-    float maxval = s_max;
-
-    float partial_sum = 0.0f;
-    for (int i = threadIdx.x; i <= own_pos; i += BLOCK_DIM) {
-        float ev = expf(inv_temperature * (a[i] - maxval));
-        o[i] = ev;
-        partial_sum += ev;
-    }
-
-    partial_sum = warpReduceSum(partial_sum);
-
-    __shared__ float partial_sums_s[BLOCK_DIM / WARP_SIZE];
-    if (threadIdx.x % WARP_SIZE == 0) {
-        partial_sums_s[threadIdx.x / WARP_SIZE] = partial_sum;
-    }
-    __syncthreads();
-
-    if (threadIdx.x < WARP_SIZE) {
-        partial_sum = (threadIdx.x < BLOCK_DIM / WARP_SIZE) ? partial_sums_s[threadIdx.x] : 0.0f;
-        partial_sum = warpReduceSum(partial_sum);
-    }
-
-    __shared__ float s_sum;
-    if (threadIdx.x == 0) {
-        s_sum = partial_sum;
-    }
-    __syncthreads();
-    float norm = 1.0f / s_sum;
-
-    for (int i = threadIdx.x; i <= own_pos; i += BLOCK_DIM) {
-        o[i] *= norm;
-    }
-}
-
-#endif // __SOFTMAX_KERNEL_CUH__

kernels_req_5/local_attention.cuh

@@ -1,9 +1,6 @@
 #ifndef __ATTENTION_CUH__
 #define __ATTENTION_CUH__
-
+#define WINDOW 128
-// WINDOW = 129 because the CPU reference uses an inclusive range [q_t - 128, q_t],
-// which covers 129 positions (128 look-back steps + the current token itself).
-#define WINDOW 129
 
 #include <cuda_runtime.h>
 #include "../utils/cuda_utils.cuh"

kernels_req_5/local_softmax.cuh

@@ -1,5 +1,6 @@
 #ifndef __SOFTMAX_KERNEL_CUH__
 #define __SOFTMAX_KERNEL_CUH__
+#define WINDOW 129
 
 #include <cuda_runtime.h>
 #include <math.h>

old_verify_new.slurm

@@ -22,7 +22,7 @@ module load cuda
 echo -e "job $SLURM_JOBID is starting on `hostname`\n\n"
 
 # ── nsys system-level profile ─────────────────────────────────────
-srun nsys profile --stats=true ./test_gpt2 > test_gpt2_nsys_prof_req_0.out
+srun nsys profile --stats=true ./build/test_gpt2 > test_gpt2_nsys_prof_comp.out
 
 # ── ncu kernel-level profile ──────────────────────────────────────
-# ncu --set full -f -o analysis_file ./test_gpt2 > test_gpt2_ncu.out
+# ncu --set full -f -o analysis_file ./build/test_gpt2 > test_gpt2_ncu.out

test_gpt2_nsys_prof_req_2.out

@@ -4,111 +4,103 @@ Batch Size: 4
 Sequence Length: 64
 GPT2 - Allocated 221 MiB for activations
 test-gpt2 - Printing 10 Expected vs Actual Logits: 
--43.431618, -43.431801
+-43.431618, -43.363079
--39.836346, -39.836525
+-39.836346, -39.774330
--43.065910, -43.066086
+-43.065910, -43.004066
--42.828045, -42.828217
+-42.828045, -42.763325
--43.529541, -43.529724
+-43.529541, -43.461735
--44.318398, -44.318581
+-44.318398, -44.251251
--41.227425, -41.227596
+-41.227425, -41.161076
--41.270760, -41.270927
+-41.270760, -41.205547
--42.541393, -42.541595
+-42.541393, -42.472733
--42.394997, -42.395184
+-42.394997, -42.328682
----------------- Logits OK! (max err 0.000069, RMSE 0.000002) ----------------
+---------------- Logits OK! (max err 0.036241, RMSE 0.001194) ----------------
 _________________________ All tests passed! _________________________
 Collecting data...
-Generating '/tmp/nsys-report-e25f.qdstrm'
+Generating '/tmp/nsys-report-c9b3.qdstrm'
-␍[1/8] [0%                          ] report7.nsys-rep␍[1/8] [0%                          ] report7.nsys-rep␍[1/8] [=16%                        ] report7.nsys-rep␍[1/8] [=====31%                    ] report7.nsys-rep␍[1/8] [=====32%                    ] report7.nsys-rep␍[1/8] [====25%                     ] report7.nsys-rep␍[1/8] [======35%                   ] report7.nsys-rep␍[1/8] [=======36%                  ] report7.nsys-rep␍[1/8] [=======37%                  ] report7.nsys-rep␍[1/8] [=======38%                  ] report7.nsys-rep␍[1/8] [=======39%                  ] report7.nsys-rep␍[1/8] [========40%                 ] report7.nsys-rep␍[1/8] [========41%                 ] report7.nsys-rep␍[1/8] [========42%                 ] report7.nsys-rep␍[1/8] [==================78%       ] report7.nsys-rep␍[1/8] [=====================86%    ] report7.nsys-rep␍[1/8] [========================100%] report7.nsys-rep␍[1/8] [========================100%] report7.nsys-rep
+␍[1/8] [0%                          ] report8.nsys-rep␍[1/8] [0%                          ] report8.nsys-rep␍[1/8] [====28%                     ] report8.nsys-rep␍[1/8] [=====29%                    ] report8.nsys-rep␍[1/8] [=====30%                    ] report8.nsys-rep␍[1/8] [=======39%                  ] report8.nsys-rep␍[1/8] [==========47%               ] report8.nsys-rep␍[1/8] [==========48%               ] report8.nsys-rep␍[1/8] [==========49%               ] report8.nsys-rep␍[1/8] [===========50%              ] report8.nsys-rep␍[1/8] [===========51%              ] report8.nsys-rep␍[1/8] [============57%             ] report8.nsys-rep␍[1/8] [=============58%            ] report8.nsys-rep␍[1/8] [=============59%            ] report8.nsys-rep␍[1/8] [=============60%            ] report8.nsys-rep␍[1/8] [==============61%           ] report8.nsys-rep␍[1/8] [==============62%           ] report8.nsys-rep␍[1/8] [==============63%           ] report8.nsys-rep␍[1/8] [=================74%        ] report8.nsys-rep␍[1/8] [=======================94%  ] report8.nsys-rep␍[1/8] [========================100%] report8.nsys-rep␍[1/8] [========================100%] report8.nsys-rep
-
[2/8] [0%                          ] report7.sqlite
[2/8] [1%                          ] report7.sqlite
[2/8] [2%                          ] report7.sqlite
[2/8] [3%                          ] report7.sqlite
[2/8] [4%                          ] report7.sqlite
[2/8] [5%                          ] report7.sqlite
[2/8] [6%                          ] report7.sqlite
[2/8] [7%                          ] report7.sqlite
[2/8] [8%                          ] report7.sqlite
[2/8] [9%                          ] report7.sqlite
[2/8] [10%                         ] report7.sqlite
[2/8] [11%                         ] report7.sqlite
[2/8] [12%                         ] report7.sqlite
[2/8] [13%                         ] report7.sqlite
[2/8] [14%                         ] report7.sqlite
[2/8] [=15%                        ] report7.sqlite
[2/8] [=16%                        ] report7.sqlite
[2/8] [=17%                        ] report7.sqlite
[2/8] [==18%                       ] report7.sqlite
[2/8] [==19%                       ] report7.sqlite
[2/8] [==20%                       ] report7.sqlite
[2/8] [==21%                       ] report7.sqlite
[2/8] [===22%                      ] report7.sqlite
[2/8] [===23%                      ] report7.sqlite
[2/8] [===24%                      ] report7.sqlite
[2/8] [====25%                     ] report7.sqlite
[2/8] [====26%                     ] report7.sqlite
[2/8] [====27%                     ] report7.sqlite
[2/8] [====28%                     ] report7.sqlite
[2/8] [=====29%                    ] report7.sqlite
[2/8] [=====30%                    ] report7.sqlite
[2/8] [=====31%                    ] report7.sqlite
[2/8] [=====32%                    ] report7.sqlite
[2/8] [======33%                   ] report7.sqlite
[2/8] [======34%                   ] report7.sqlite
[2/8] [======35%                   ] report7.sqlite
[2/8] [=======36%                  ] report7.sqlite
[2/8] [=======37%                  ] report7.sqlite
[2/8] [=======38%                  ] report7.sqlite
[2/8] [=======39%                  ] report7.sqlite
[2/8] [========40%                 ] report7.sqlite
[2/8] [========41%                 ] report7.sqlite
[2/8] [========42%                 ] report7.sqlite
[2/8] [=========43%                ] report7.sqlite
[2/8] [=========44%                ] report7.sqlite
[2/8] [=========45%                ] report7.sqlite
[2/8] [=========46%                ] report7.sqlite
[2/8] [==========47%               ] report7.sqlite
[2/8] [==========48%               ] report7.sqlite
[2/8] [==========49%               ] report7.sqlite
[2/8] [===========50%              ] report7.sqlite
[2/8] [===========51%              ] report7.sqlite
[2/8] [===========52%              ] report7.sqlite
[2/8] [===========53%              ] report7.sqlite
[2/8] [============54%             ] report7.sqlite
[2/8] [============55%             ] report7.sqlite
[2/8] [============56%             ] report7.sqlite
[2/8] [============57%             ] report7.sqlite
[2/8] [=============58%            ] report7.sqlite
[2/8] [=============59%            ] report7.sqlite
[2/8] [=============60%            ] report7.sqlite
[2/8] [==============61%           ] report7.sqlite
[2/8] [==============62%           ] report7.sqlite
[2/8] [==============63%           ] report7.sqlite
[2/8] [==============64%           ] report7.sqlite
[2/8] [===============65%          ] report7.sqlite
[2/8] [===============66%          ] report7.sqlite
[2/8] [===============67%          ] report7.sqlite
[2/8] [================68%         ] report7.sqlite
[2/8] [================69%         ] report7.sqlite
[2/8] [================70%         ] report7.sqlite
[2/8] [================71%         ] report7.sqlite
[2/8] [=================72%        ] report7.sqlite
[2/8] [=================73%        ] report7.sqlite
[2/8] [=================74%        ] report7.sqlite
[2/8] [==================75%       ] report7.sqlite
[2/8] [==================76%       ] report7.sqlite
[2/8] [==================77%       ] report7.sqlite
[2/8] [==================78%       ] report7.sqlite
[2/8] [===================79%      ] report7.sqlite
[2/8] [===================80%      ] report7.sqlite
[2/8] [===================81%      ] report7.sqlite
[2/8] [===================82%      ] report7.sqlite
[2/8] [====================83%     ] report7.sqlite
[2/8] [====================84%     ] report7.sqlite
[2/8] [====================85%     ] report7.sqlite
[2/8] [=====================86%    ] report7.sqlite
[2/8] [=====================87%    ] report7.sqlite
[2/8] [=====================88%    ] report7.sqlite
[2/8] [=====================89%    ] report7.sqlite
[2/8] [======================90%   ] report7.sqlite
[2/8] [======================91%   ] report7.sqlite
[2/8] [======================92%   ] report7.sqlite
[2/8] [=======================93%  ] report7.sqlite
[2/8] [=======================94%  ] report7.sqlite
[2/8] [=======================95%  ] report7.sqlite
[2/8] Line truncated
+
[2/8] [0%                          ] report8.sqlite
[2/8] [1%                          ] report8.sqlite
[2/8] [2%                          ] report8.sqlite
[2/8] [3%                          ] report8.sqlite
[2/8] [4%                          ] report8.sqlite
[2/8] [5%                          ] report8.sqlite
[2/8] [6%                          ] report8.sqlite
[2/8] [7%                          ] report8.sqlite
[2/8] [8%                          ] report8.sqlite
[2/8] [9%                          ] report8.sqlite
[2/8] [10%                         ] report8.sqlite
[2/8] [11%                         ] report8.sqlite
[2/8] [12%                         ] report8.sqlite
[2/8] [13%                         ] report8.sqlite
[2/8] [14%                         ] report8.sqlite
[2/8] [=15%                        ] report8.sqlite
[2/8] [=16%                        ] report8.sqlite
[2/8] [=17%                        ] report8.sqlite
[2/8] [==18%                       ] report8.sqlite
[2/8] [==19%                       ] report8.sqlite
[2/8] [==20%                       ] report8.sqlite
[2/8] [==21%                       ] report8.sqlite
[2/8] [===22%                      ] report8.sqlite
[2/8] [===23%                      ] report8.sqlite
[2/8] [===24%                      ] report8.sqlite
[2/8] [====25%                     ] report8.sqlite
[2/8] [====26%                     ] report8.sqlite
[2/8] [====27%                     ] report8.sqlite
[2/8] [====28%                     ] report8.sqlite
[2/8] [=====29%                    ] report8.sqlite
[2/8] [=====30%                    ] report8.sqlite
[2/8] [=====31%                    ] report8.sqlite
[2/8] [=====32%                    ] report8.sqlite
[2/8] [======33%                   ] report8.sqlite
[2/8] [======34%                   ] report8.sqlite
[2/8] [======35%                   ] report8.sqlite
[2/8] [=======36%                  ] report8.sqlite
[2/8] [=======37%                  ] report8.sqlite
[2/8] [=======38%                  ] report8.sqlite
[2/8] [=======39%                  ] report8.sqlite
[2/8] [========40%                 ] report8.sqlite
[2/8] [========41%                 ] report8.sqlite
[2/8] [========42%                 ] report8.sqlite
[2/8] [=========43%                ] report8.sqlite
[2/8] [=========44%                ] report8.sqlite
[2/8] [=========45%                ] report8.sqlite
[2/8] [=========46%                ] report8.sqlite
[2/8] [==========47%               ] report8.sqlite
[2/8] [==========48%               ] report8.sqlite
[2/8] [==========49%               ] report8.sqlite
[2/8] [===========50%              ] report8.sqlite
[2/8] [===========51%              ] report8.sqlite
[2/8] [===========52%              ] report8.sqlite
[2/8] [===========53%              ] report8.sqlite
[2/8] [============54%             ] report8.sqlite
[2/8] [============55%             ] report8.sqlite
[2/8] [============56%             ] report8.sqlite
[2/8] [============57%             ] report8.sqlite
[2/8] [=============58%            ] report8.sqlite
[2/8] [=============59%            ] report8.sqlite
[2/8] [=============60%            ] report8.sqlite
[2/8] [==============61%           ] report8.sqlite
[2/8] [==============62%           ] report8.sqlite
[2/8] [==============63%           ] report8.sqlite
[2/8] [==============64%           ] report8.sqlite
[2/8] [===============65%          ] report8.sqlite
[2/8] [===============66%          ] report8.sqlite
[2/8] [===============67%          ] report8.sqlite
[2/8] [================68%         ] report8.sqlite
[2/8] [================69%         ] report8.sqlite
[2/8] [================70%         ] report8.sqlite
[2/8] [================71%         ] report8.sqlite
[2/8] [=================72%        ] report8.sqlite
[2/8] [=================73%        ] report8.sqlite
[2/8] [=================74%        ] report8.sqlite
[2/8] [==================75%       ] report8.sqlite
[2/8] [==================76%       ] report8.sqlite
[2/8] [==================77%       ] report8.sqlite
[2/8] [==================78%       ] report8.sqlite
[2/8] [===================79%      ] report8.sqlite
[2/8] [===================80%      ] report8.sqlite
[2/8] [===================81%      ] report8.sqlite
[2/8] [===================82%      ] report8.sqlite
[2/8] [====================83%     ] report8.sqlite
[2/8] [====================84%     ] report8.sqlite
[2/8] [====================85%     ] report8.sqlite
[2/8] [=====================86%    ] report8.sqlite
[2/8] [=====================87%    ] report8.sqlite
[2/8] [=====================88%    ] report8.sqlite
[2/8] [=====================89%    ] report8.sqlite
[2/8] [======================90%   ] report8.sqlite
[2/8] [======================91%   ] report8.sqlite
[2/8] [======================92%   ] report8.sqlite
[2/8] [=======================93%  ] report8.sqlite
[2/8] [=======================94%  ] report8.sqlite
[2/8] [=======================95%  ] report8.sqlite
[2/8] Line truncated
 [3/8] Executing 'nvtx_sum' stats report
 [4/8] Executing 'osrt_sum' stats report
 
- Time (%)  Total Time (ns)  Num Calls   Avg (ns)    Med (ns)   Min (ns)  Max (ns)   StdDev (ns)           Name         
+ Time (%)  Total Time (ns)  Num Calls   Avg (ns)     Med (ns)    Min (ns)  Max (ns)   StdDev (ns)          Name         
- --------  ---------------  ---------  ----------  ----------  --------  ---------  -----------  ----------------------
+ --------  ---------------  ---------  -----------  -----------  --------  ---------  -----------  ---------------------
-     62.0        680717261         17  40042191.8  10065417.0      1532  100136863   44760103.1  poll                  
+     54.7       1713974843         27   63480549.7  100113755.0      1193  108696839   46506883.0  poll                 
-     14.8        162110818         38   4266074.2      6131.0      3496  161029552   26117864.7  fopen                 
+     36.4       1138051148          6  189675191.3   15085903.5      7013  820460548  328760934.6  fread                
-     11.0        120373004          6  20062167.3     89457.5      6833   95159843   38120924.5  fread                 
+      4.9        153194333        595     257469.5      22662.0      1032   18701261     996184.4  ioctl                
-     10.4        114685549        591    194053.4     23404.0      1032   22556281     980200.7  ioctl                 
+      3.8        118955136         38    3130398.3       6247.0      3576  118016722   19140872.4  fopen                
-      0.8          9116342          1   9116342.0   9116342.0   9116342    9116342          0.0  pthread_cond_broadcast
+      0.1          2406150         25      96246.0       8916.0      6703    1662941     329260.7  mmap64               
-      0.4          3892842         32    121651.3      2144.0      1353    3504445     619781.4  fclose                
+      0.1          1648944         10     164894.4      30456.5     21350     809013     285878.4  write                
-      0.2          2358791         25     94351.6      8787.0      6392    1637323     324262.9  mmap64                
+      0.0           781242         32      24413.8       2209.0      1333     381844      85901.8  fclose               
-      0.2          1996705          1   1996705.0   1996705.0   1996705    1996705          0.0  pthread_cond_wait     
+      0.0           578514         17      34030.2       4980.0      2104     377536      90206.1  mmap                 
-      0.2          1950609         10    195060.9     27181.0     22051     798313     285177.5  write                 
+      0.0           413214          2     206607.0     206607.0    163236     249978      61335.9  pthread_create       
-      0.1           591677         17     34804.5      5981.0      1893     395630      94748.6  mmap                  
+      0.0           322402         55       5861.9       5991.0      2354      11190       2184.1  open64               
-      0.0           362288          2    181144.0    181144.0    171561     190727      13552.4  pthread_create        
+      0.0            76614          1      76614.0      76614.0     76614      76614          0.0  fgets                
-      0.0           322858         55      5870.1      5861.0      2314      14327       2395.2  open64                
+      0.0            61595          1      61595.0      61595.0     61595      61595          0.0  fflush               
-      0.0            76734          1     76734.0     76734.0     76734      76734          0.0  fgets                 
+      0.0            43209          8       5401.1       4368.0      2996      11862       2985.1  munmap               
-      0.0            62457          1     62457.0     62457.0     62457      62457          0.0  fflush                
+      0.0            38103          6       6350.5       6252.0      3076      10069       2427.2  open                 
-      0.0            40205          6      6700.8      6922.5      3527       9608       2269.3  open                  
+      0.0            24296          3       8098.7       7855.0      3567      12874       4658.3  pipe2                
-      0.0            38922          8      4865.3      5074.5      2765       7153       1707.3  munmap                
+      0.0            18094          9       2010.4       1663.0      1032       3627        895.0  read                 
-      0.0            19787          3      6595.7      7844.0      3657       8286       2554.5  pipe2                 
+      0.0            16701          2       8350.5       8350.5      5090      11611       4611.0  socket               
-      0.0            18904         10      1890.4      1623.0      1012       3516        882.0  read                  
+      0.0            11391          1      11391.0      11391.0     11391      11391          0.0  connect              
-      0.0            15960          2      7980.0      7980.0      5100      10860       4072.9  socket                
+      0.0             6843          1       6843.0       6843.0      6843       6843          0.0  fopen64              
-      0.0            11101          1     11101.0     11101.0     11101      11101          0.0  connect               
+      0.0             3707          2       1853.5       1853.5      1453       2254        566.4  fwrite               
-      0.0             6472          1      6472.0      6472.0      6472       6472          0.0  fopen64               
+      0.0             2465          1       2465.0       2465.0      2465       2465          0.0  bind                 
-      0.0             3417          2      1708.5      1708.5      1393       2024        446.2  fwrite                
+      0.0             2114          2       1057.0       1057.0      1042       1072         21.2  fcntl                
-      0.0             1593          1      1593.0      1593.0      1593       1593          0.0  bind                  
+      0.0             1062          1       1062.0       1062.0      1062       1062          0.0  listen               
+      0.0             1042          1       1042.0       1042.0      1042       1042          0.0  pthread_mutex_trylock
 
 [5/8] Executing 'cuda_api_sum' stats report
 
  Time (%)  Total Time (ns)  Num Calls   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)                Name               
  --------  ---------------  ---------  ----------  ----------  --------  --------  -----------  ---------------------------------
-     61.4         51907511          3  17302503.7  15637119.0     21119  36249273   18171403.8  cudaMemcpy                       
+     75.3         55764098          3  18588032.7  18908389.0     17473  36838236   18412471.8  cudaMemcpy                       
-     18.9         15960434          4   3990108.5   3249843.5   1031059   8429688    3147630.5  cuLibraryLoadData                
+     13.8         10194027         12    849502.3     11000.0       180   5616656    1649151.5  cudaFree                         
-      6.1          5177442         12    431453.5      6161.5       191   1837587     621479.2  cudaFree                         
+      3.3          2407413          1   2407413.0   2407413.0   2407413   2407413          0.0  cudaFreeHost                     
-      5.7          4819182       1347      3577.7      3086.0      2835    335788       9243.6  cudaLaunchKernel                 
+      2.2          1633015          1   1633015.0   1633015.0   1633015   1633015          0.0  cudaGetDeviceProperties_v2_v12000
-      2.9          2489386          1   2489386.0   2489386.0   2489386   2489386          0.0  cudaFreeHost                     
+      1.8          1328725          1   1328725.0   1328725.0   1328725   1328725          0.0  cudaMallocHost                   
-      1.7          1448309          1   1448309.0   1448309.0   1448309   1448309          0.0  cudaGetDeviceProperties_v2_v12000
+      1.5          1095220        171      6404.8      3026.0      2846    406250      31186.8  cudaLaunchKernel                 
-      1.6          1383518          1   1383518.0   1383518.0   1383518   1383518          0.0  cudaMallocHost                   
+      1.1           837937          6    139656.2    131245.5      4819    364852     133988.9  cudaMalloc                       
-      0.9           766724          6    127787.3    123215.5      3447    318305     118859.5  cudaMalloc                       
+      0.6           426397         18     23688.7       295.0       260    416068      97927.8  cudaEventCreateWithFlags         
-      0.2           160077        838       191.0       151.0        70      6683        242.2  cuGetProcAddress_v2              
+      0.3           225714        838       269.3       180.0        90     44584       1538.5  cuGetProcAddress_v2              
-      0.2           139420         48      2904.6      2500.0      2184     13164       1581.6  cudaMemsetAsync                  
+      0.1            59281          1     59281.0     59281.0     59281     59281          0.0  cudaMemset                       
-      0.2           137165         48      2857.6      2690.0      1974     12383       1519.2  cudaEventRecord                  
+      0.0            21260          4      5315.0      3978.0      1162     12142       4779.4  cudaDeviceSynchronize            
-      0.0            35015          1     35015.0     35015.0     35015     35015          0.0  cudaMemset                       
+      0.0             9335         18       518.6       290.5       210      3466        760.8  cudaEventDestroy                 
-      0.0            30347          4      7586.8      7734.0      6913      7966        469.4  cuLibraryGetKernel               
+      0.0             5982          4      1495.5      1513.0      1293      1663        172.9  cuInit                           
-      0.0            27345        240       113.9        90.0        69      1262         81.3  cuStreamGetCaptureInfo_v2        
+      0.0             1283          2       641.5       641.5       261      1022        538.1  cudaGetDriverEntryPoint_v11030   
-      0.0            17219         18       956.6       300.0       260      7834       1819.2  cudaEventCreateWithFlags         
+      0.0             1233          3       411.0       131.0        90      1012        520.9  cuModuleGetLoadingMode           
-      0.0            12852          4      3213.0      2549.5       991      6762       2485.5  cudaDeviceSynchronize            
-      0.0             8116         18       450.9       301.0       211      2675        566.0  cudaEventDestroy                 
-      0.0             5371          4      1342.8      1343.0      1172      1513        139.5  cuInit                           
-      0.0             3317          1      3317.0      3317.0      3317      3317          0.0  cudaEventQuery                   
-      0.0             1322          3       440.7       180.0       130      1012        495.4  cuModuleGetLoadingMode           
-      0.0              872          2       436.0       436.0       211       661        318.2  cudaGetDriverEntryPoint_v11030   
 
 [6/8] Executing 'cuda_gpu_kern_sum' stats report
 
  Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                                                  Name                                                
  --------  ---------------  ---------  --------  --------  --------  --------  -----------  ----------------------------------------------------------------------------------------------------
-     25.6          4083613         25  163344.5  163744.0    158464    166208       2004.3  layernorm_forward_kernel(float *, float *, float *, const float *, const float *, const float *, in…
+     90.7         10611176         49  216554.6  162239.0     77280   2391507     328408.1  matmul_forward_kernel(float *, const float *, const float *, const float *, int, int, int, int)     
-     22.8          3644222        576    6326.8    6336.0      6144      6592         79.1  ampere_sgemm_32x128_nn                                                                              
+      1.9           226655         12   18887.9   18880.0     18784     18944         41.2  gelu_forward_kernel(float *, const float *, int)                                                    
-     16.8          2675550         25  107022.0   78560.0     60384   1000383     186333.0  ampere_sgemm_128x64_tn                                                                              
+      1.2           145759         24    6073.3    6080.0      5663      6528        378.6  residual_forward_kernel(float *, float *, float *, int)                                             
-     15.9          2535902        576    4402.6    4384.0      4352      4864         60.7  ampere_sgemm_32x32_sliced1x4_tn                                                                     
+      1.2           141087         12   11757.3   11728.0     11648     12000         97.8  softmax_forward_kernel(float *, float, const float *, int, int)                                     
-     10.0          1591582         12  132631.8  133168.0    129568    134783       1430.5  ampere_sgemm_64x64_tn                                                                               
+      1.1           131199         25    5248.0    5088.0      4672      5984        440.1  layernorm_forward_kernel(float *, float *, float *, const float *, const float *, const float *, in…
-      2.4           390976         12   32581.3   32528.0     32384     32928        170.9  ampere_sgemm_64x32_sliced1x4_tn                                                                     
+      0.9            99744         12    8312.0    8288.0      8224      8448         62.7  attention_v_kernel(float *, float *, float *, int, int, int, int)                                   
-      1.4           223072         12   18589.3   18560.0     18528     18848         90.0  gelu_forward_kernel(float *, const float *, int)                                                    
+      0.8            95968         12    7997.3    7968.0      7904      8160         72.8  preattention_scores_kernel(float *, float *, int, int, int, int, float *)                           
-      1.4           219968         48    4582.7    4288.0      2848      7616       1622.8  bias_add_kernel(float *, const float *, int, int)                                                   
+      0.8            94622         12    7885.2    7856.0      7648      8255        167.1  permute_kernel(float *, float *, float *, const float *, int, int, int, int)                        
-      1.4           217344         12   18112.0   18000.0     17920     19200        353.5  softmax_forward_kernel(float *, float, const float *, int, int)                                     
+      0.7            86176          1   86176.0   86176.0     86176     86176          0.0  encoder_forward_kernel(float *, const int *, const float *, const float *, int, int, int)           
-      0.9           142080         24    5920.0    5952.0      5632      6240        265.2  residual_forward_kernel(float *, float *, float *, int)                                             
+      0.6            70623         12    5885.3    5888.0      5856      5920         21.2  unpermute_kernel(float *, float *, int, int, int, int)                                              
-      0.6            89376          1   89376.0   89376.0     89376     89376          0.0  encoder_forward_kernel(float *, const int *, const float *, const float *, int, int, int)           
-      0.5            84511         12    7042.6    7008.0      6976      7392        112.9  permute_kernel(float *, float *, float *, const float *, int, int, int, int)                        
-      0.4            70944         12    5912.0    5888.0      5824      6016         62.7  unpermute_kernel(float *, float *, int, int, int, int)                                              
 
 [7/8] Executing 'cuda_gpu_mem_time_sum' stats report
 
  Time (%)  Total Time (ns)  Count   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)           Operation          
  --------  ---------------  -----  ----------  ----------  --------  --------  -----------  ----------------------------
-     81.3         36158470      2  18079235.0  18079235.0       832  36157638   25566722.7  [CUDA memcpy Host-to-Device]
+     81.9         36690814      2  18345407.0  18345407.0       864  36689950   25943101.5  [CUDA memcpy Host-to-Device]
-     17.6          7821051      1   7821051.0   7821051.0   7821051   7821051          0.0  [CUDA memcpy Device-to-Host]
+     17.2          7698135      1   7698135.0   7698135.0   7698135   7698135          0.0  [CUDA memcpy Device-to-Host]
-      1.1           473504     49      9663.3      1184.0       480    407616      58036.0  [CUDA memset]               
+      0.9           407997      1    407997.0    407997.0    407997    407997          0.0  [CUDA memset]               
 
 [8/8] Executing 'cuda_gpu_mem_size_sum' stats report
 
  Total (MB)  Count  Avg (MB)  Med (MB)  Min (MB)  Max (MB)  StdDev (MB)           Operation          
  ----------  -----  --------  --------  --------  --------  -----------  ----------------------------
     497.905      2   248.952   248.952     0.001   497.904      352.070  [CUDA memcpy Host-to-Device]
-    232.407     49     4.743     0.000     0.000   232.392       33.199  [CUDA memset]               
+    232.392      1   232.392   232.392   232.392   232.392        0.000  [CUDA memset]               
      51.511      1    51.511    51.511    51.511    51.511        0.000  [CUDA memcpy Device-to-Host]
 
 Generated:
-	/u/asingh41/sp26_ece408_Temperature1.5/report7.nsys-rep
+	/u/asingh41/sp26_ece408_Temperature1.5/report8.nsys-rep
-	/u/asingh41/sp26_ece408_Temperature1.5/report7.sqlite
+	/u/asingh41/sp26_ece408_Temperature1.5/report8.sqlite

test_gpt2_nsys_prof_req_perf.out

@@ -0,0 +1,106 @@
+test-gpt2-kernels - TF32 Mode is turned OFF.
+test-gpt2 - [State]
+Batch Size: 4
+Sequence Length: 64
+GPT2 - Allocated 221 MiB for activations
+test-gpt2 - Printing 10 Expected vs Actual Logits: 
+-43.431618, -43.363079
+-39.836346, -39.774330
+-43.065910, -43.004066
+-42.828045, -42.763325
+-43.529541, -43.461735
+-44.318398, -44.251251
+-41.227425, -41.161076
+-41.270760, -41.205547
+-42.541393, -42.472733
+-42.394997, -42.328682
+---------------- Logits OK! (max err 0.036241, RMSE 0.001194) ----------------
+_________________________ All tests passed! _________________________
+Collecting data...
+Generating '/tmp/nsys-report-8038.qdstrm'
+␍[1/8] [0%                          ] report10.nsys-rep␍[1/8] [0%                          ] report10.nsys-rep␍[1/8] [====26%                     ] report10.nsys-rep␍[1/8] [==========48%               ] report10.nsys-rep␍[1/8] [==========49%               ] report10.nsys-rep␍[1/8] [===========50%              ] report10.nsys-rep␍[1/8] [=====32%                    ] report10.nsys-rep␍[1/8] [===========50%              ] report10.nsys-rep␍[1/8] [===========51%              ] report10.nsys-rep␍[1/8] [===========52%              ] report10.nsys-rep␍[1/8] [=============58%            ] report10.nsys-rep␍[1/8] [==================75%       ] report10.nsys-rep␍[1/8] [======================91%   ] report10.nsys-rep␍[1/8] [========================100%] report10.nsys-rep␍[1/8] [========================100%] report10.nsys-rep
+
[2/8] [0%                          ] report10.sqlite
[2/8] [1%                          ] report10.sqlite
[2/8] [2%                          ] report10.sqlite
[2/8] [3%                          ] report10.sqlite
[2/8] [4%                          ] report10.sqlite
[2/8] [5%                          ] report10.sqlite
[2/8] [6%                          ] report10.sqlite
[2/8] [7%                          ] report10.sqlite
[2/8] [8%                          ] report10.sqlite
[2/8] [9%                          ] report10.sqlite
[2/8] [10%                         ] report10.sqlite
[2/8] [11%                         ] report10.sqlite
[2/8] [12%                         ] report10.sqlite
[2/8] [13%                         ] report10.sqlite
[2/8] [14%                         ] report10.sqlite
[2/8] [=15%                        ] report10.sqlite
[2/8] [=16%                        ] report10.sqlite
[2/8] [=17%                        ] report10.sqlite
[2/8] [==18%                       ] report10.sqlite
[2/8] [==19%                       ] report10.sqlite
[2/8] [==20%                       ] report10.sqlite
[2/8] [==21%                       ] report10.sqlite
[2/8] [===22%                      ] report10.sqlite
[2/8] [===23%                      ] report10.sqlite
[2/8] [===24%                      ] report10.sqlite
[2/8] [====25%                     ] report10.sqlite
[2/8] [====26%                     ] report10.sqlite
[2/8] [====27%                     ] report10.sqlite
[2/8] [====28%                     ] report10.sqlite
[2/8] [=====29%                    ] report10.sqlite
[2/8] [=====30%                    ] report10.sqlite
[2/8] [=====31%                    ] report10.sqlite
[2/8] [=====32%                    ] report10.sqlite
[2/8] [======33%                   ] report10.sqlite
[2/8] [======34%                   ] report10.sqlite
[2/8] [======35%                   ] report10.sqlite
[2/8] [=======36%                  ] report10.sqlite
[2/8] [=======37%                  ] report10.sqlite
[2/8] [=======38%                  ] report10.sqlite
[2/8] [=======39%                  ] report10.sqlite
[2/8] [========40%                 ] report10.sqlite
[2/8] [========41%                 ] report10.sqlite
[2/8] [========42%                 ] report10.sqlite
[2/8] [=========43%                ] report10.sqlite
[2/8] [=========44%                ] report10.sqlite
[2/8] [=========45%                ] report10.sqlite
[2/8] [=========46%                ] report10.sqlite
[2/8] [==========47%               ] report10.sqlite
[2/8] [==========48%               ] report10.sqlite
[2/8] [==========49%               ] report10.sqlite
[2/8] [===========50%              ] report10.sqlite
[2/8] [===========51%              ] report10.sqlite
[2/8] [===========52%              ] report10.sqlite
[2/8] [===========53%              ] report10.sqlite
[2/8] [============54%             ] report10.sqlite
[2/8] [============55%             ] report10.sqlite
[2/8] [============56%             ] report10.sqlite
[2/8] [============57%             ] report10.sqlite
[2/8] [=============58%            ] report10.sqlite
[2/8] [=============59%            ] report10.sqlite
[2/8] [=============60%            ] report10.sqlite
[2/8] [==============61%           ] report10.sqlite
[2/8] [==============62%           ] report10.sqlite
[2/8] [==============63%           ] report10.sqlite
[2/8] [==============64%           ] report10.sqlite
[2/8] [===============65%          ] report10.sqlite
[2/8] [===============66%          ] report10.sqlite
[2/8] [===============67%          ] report10.sqlite
[2/8] [================68%         ] report10.sqlite
[2/8] [================69%         ] report10.sqlite
[2/8] [================70%         ] report10.sqlite
[2/8] [================71%         ] report10.sqlite
[2/8] [=================72%        ] report10.sqlite
[2/8] [=================73%        ] report10.sqlite
[2/8] [=================74%        ] report10.sqlite
[2/8] [==================75%       ] report10.sqlite
[2/8] [==================76%       ] report10.sqlite
[2/8] [==================77%       ] report10.sqlite
[2/8] [==================78%       ] report10.sqlite
[2/8] [===================79%      ] report10.sqlite
[2/8] [===================80%      ] report10.sqlite
[2/8] [===================81%      ] report10.sqlite
[2/8] [===================82%      ] report10.sqlite
[2/8] [====================83%     ] report10.sqlite
[2/8] [====================84%     ] report10.sqlite
[2/8] [====================85%     ] report10.sqlite
[2/8] [=====================86%    ] report10.sqlite
[2/8] [=====================87%    ] report10.sqlite
[2/8] [=====================88%    ] report10.sqlite
[2/8] [=====================89%    ] report10.sqlite
[2/8] [======================90%   ] report10.sqlite
[2/8] [======================91%   ] report10.sqlite
[2/8] [======================92%   ] report10.sqlite
[2/8] [=======================93%  ] report10.sqlite
[2/8] [=========Line truncated
+[3/8] Executing 'nvtx_sum' stats report
+[4/8] Executing 'osrt_sum' stats report
+
+ Time (%)  Total Time (ns)  Num Calls   Avg (ns)     Med (ns)   Min (ns)  Max (ns)   StdDev (ns)           Name         
+ --------  ---------------  ---------  -----------  ----------  --------  ---------  -----------  ----------------------
+     56.0        996338383         20   49816919.1  37144026.0      1233  100147022   47273324.4  poll                  
+     34.2        608280586          6  101380097.7    818540.0      7524  540651394  216790051.0  fread                 
+      7.5        133295089        592     225160.6     22788.0      1002   20761770     945697.8  ioctl                 
+      1.3         22334199         38     587742.1      6277.0      3637   21155518    3428230.9  fopen                 
+      0.4          7061795          1    7061795.0   7061795.0   7061795    7061795          0.0  pthread_cond_broadcast
+      0.4          6959863          1    6959863.0   6959863.0   6959863    6959863          0.0  pthread_cond_wait     
+      0.1          2388902         25      95556.1      9268.0      6172    1630658     322960.7  mmap64                
+      0.1          1595422         10     159542.2     32852.0     22562     552156     214109.1  write                 
+      0.0           738626         32      23082.1      2129.0      1352     379712      82280.0  fclose                
+      0.0           572887         17      33699.2      5110.0      1984     378079      90516.0  mmap                  
+      0.0           359054          2     179527.0    179527.0    164950     194104      20615.0  pthread_create        
+      0.0           331892         55       6034.4      5961.0      2354      12935       2361.7  open64                
+      0.0            76394          1      76394.0     76394.0     76394      76394          0.0  fgets                 
+      0.0            65062          1      65062.0     65062.0     65062      65062          0.0  fflush                
+      0.0            56836          9       6315.1      4839.0      2995      18494       4843.8  munmap                
+      0.0            36999          6       6166.5      6011.5      3125       9508       2180.5  open                  
+      0.0            20970          3       6990.0      8316.0      3286       9368       3250.6  pipe2                 
+      0.0            18394         10       1839.4      1353.0      1042       3456        943.7  read                  
+      0.0            16671          2       8335.5      8335.5      5571      11100       3909.6  socket                
+      0.0             8135          1       8135.0      8135.0      8135       8135          0.0  connect               
+      0.0             6442          1       6442.0      6442.0      6442       6442          0.0  fopen64               
+      0.0             3367          2       1683.5      1683.5      1423       1944        368.4  fwrite                
+      0.0             1984          1       1984.0      1984.0      1984       1984          0.0  bind                  
+      0.0             1182          1       1182.0      1182.0      1182       1182          0.0  listen                
+
+[5/8] Executing 'cuda_api_sum' stats report
+
+ Time (%)  Total Time (ns)  Num Calls   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)                Name               
+ --------  ---------------  ---------  ----------  ----------  --------  --------  -----------  ---------------------------------
+     80.8         58664668          3  19554889.3  21881269.0     17532  36765867   18484292.3  cudaMemcpy                       
+      8.6          6233820         12    519485.0      7173.0       190   2255630     790172.6  cudaFree                         
+      3.4          2475383          1   2475383.0   2475383.0   2475383   2475383          0.0  cudaFreeHost                     
+      2.5          1802090          1   1802090.0   1802090.0   1802090   1802090          0.0  cudaGetDeviceProperties_v2_v12000
+      1.8          1334142          1   1334142.0   1334142.0   1334142   1334142          0.0  cudaMallocHost                   
+      1.4          1040310        171      6083.7      2826.0      2675    390603      30047.0  cudaLaunchKernel                 
+      1.1           769711          6    128285.2    124803.5      5030    312856     117219.1  cudaMalloc                       
+      0.2           153352        838       183.0       150.0        90       762         90.7  cuGetProcAddress_v2              
+      0.1            41117          1     41117.0     41117.0     41117     41117          0.0  cudaMemset                       
+      0.0            20561         18      1142.3       270.5       250      5761       1731.3  cudaEventCreateWithFlags         
+      0.0            13585          4      3396.3      2785.0      1002      7013       2566.9  cudaDeviceSynchronize            
+      0.0             8154         18       453.0       296.0       230      2504        527.8  cudaEventDestroy                 
+      0.0             5150          4      1287.5      1227.5      1082      1613        233.6  cuInit                           
+      0.0             1423          3       474.3       170.0        91      1162        596.8  cuModuleGetLoadingMode           
+      0.0              801          2       400.5       400.5       260       541        198.7  cudaGetDriverEntryPoint_v11030   
+
+[6/8] Executing 'cuda_gpu_kern_sum' stats report
+
+ Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                                                  Name                                                
+ --------  ---------------  ---------  --------  --------  --------  --------  -----------  ----------------------------------------------------------------------------------------------------
+     90.6         13694036         49  279470.1  217344.0     97088   3138461     429723.5  matmul_forward_kernel(float *, const float *, const float *, const float *, int, int, int, int)     
+      2.0           302144         12   25178.7   25184.0     24960     25248         72.0  gelu_forward_kernel(float *, const float *, int)                                                    
+      1.3           189280         24    7886.7    7888.0      7488      8288        357.1  residual_forward_kernel(float *, float *, float *, int)                                             
+      1.2           186336         12   15528.0   15536.0     15392     15744        101.3  softmax_forward_kernel(float *, float, const float *, int, int)                                     
+      1.1           170272         25    6810.9    6528.0      6112      7904        602.7  layernorm_forward_kernel(float *, float *, float *, const float *, const float *, const float *, in…
+      0.8           126272         12   10522.7   10528.0     10400     10656         75.8  preattention_scores_kernel(float *, float *, int, int, int, int, float *)                           
+      0.8           125728         12   10477.3   10464.0     10368     10688         98.8  attention_v_kernel(float *, float *, float *, int, int, int, int)                                   
+      0.8           119232         12    9936.0    9920.0      9760     10112        144.7  permute_kernel(float *, float *, float *, const float *, int, int, int, int)                        
+      0.7           113183          1  113183.0  113183.0    113183    113183          0.0  encoder_forward_kernel(float *, const int *, const float *, const float *, int, int, int)           
+      0.6            94496         12    7874.7    7872.0      7840      7936         21.4  unpermute_kernel(float *, float *, int, int, int, int)                                              
+
+[7/8] Executing 'cuda_gpu_mem_time_sum' stats report
+
+ Time (%)  Total Time (ns)  Count   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)           Operation          
+ --------  ---------------  -----  ----------  ----------  --------  --------  -----------  ----------------------------
+     82.9         36637635      2  18318817.5  18318817.5      1088  36636547   25905181.5  [CUDA memcpy Host-to-Device]
+     16.2          7173307      1   7173307.0   7173307.0   7173307   7173307          0.0  [CUDA memcpy Device-to-Host]
+      0.9           409248      1    409248.0    409248.0    409248    409248          0.0  [CUDA memset]               
+
+[8/8] Executing 'cuda_gpu_mem_size_sum' stats report
+
+ Total (MB)  Count  Avg (MB)  Med (MB)  Min (MB)  Max (MB)  StdDev (MB)           Operation          
+ ----------  -----  --------  --------  --------  --------  -----------  ----------------------------
+    497.905      2   248.952   248.952     0.001   497.904      352.070  [CUDA memcpy Host-to-Device]
+    232.392      1   232.392   232.392   232.392   232.392        0.000  [CUDA memset]               
+     51.511      1    51.511    51.511    51.511    51.511        0.000  [CUDA memcpy Device-to-Host]
+
+Generated:
+	/u/asingh41/sp26_ece408_Temperature1.5/report10.nsys-rep
+	/u/asingh41/sp26_ece408_Temperature1.5/report10.sqlite

test_req0.sh

@@ -1,7 +1,22 @@
 cp kernels_req_0/attention.cuh kernels/attention.cuh 
 cp kernels_req_0/matmul.cuh kernels/matmul.cuh 
 
+sleep 3
+
+echo "changed !!!"
+
 make test_gpt2
 
+echo "building binary !!!"
+
+sleep 3
+
 git checkout kernels/attention.cuh
-git checkout kernels/matmul.cuh
+git checkout kernels/matmul.cuh
+
+echo "RUNNING !!!"
+sbatch verify.slurm
+
+squeue -u $USER
+
+sleep 3

test_req1.sh

@@ -8,4 +8,9 @@ make test_gpt2
 
 
 git checkout kernels/attention.cuh
-git checkout kernels/matmul.cuh
+git checkout kernels/matmul.cuh
+
+echo "RUNNING !!!"
+sbatch verify.slurm
+
+squeue -u $USER

test_req2.sh

@@ -1,11 +1,12 @@
 cp kernels_req_2/attention.cuh kernels/attention.cuh 
 cp kernels_req_2/matmul.cuh kernels/matmul.cuh 
 
-
-
 make test_gpt2
 
-
-
 git checkout kernels/attention.cuh
-git checkout kernels/matmul.cuh
+git checkout kernels/matmul.cuh
+
+echo "RUNNING !!!"
+sbatch verify.slurm
+
+squeue -u $USER

test_req3.sh

@@ -6,4 +6,9 @@ make test_gpt2
 
 git checkout kernels/softmax.cuh
 git checkout kernels/layernorm.cuh
-git checkout kernels/attention.cuh
+git checkout kernels/attention.cuh
+
+echo "RUNNING !!!"
+sbatch verify.slurm
+
+squeue -u $USER

verify_new.slurm

@@ -11,8 +11,8 @@
 #SBATCH --gpus-per-node=1
 #SBATCH --gpu-bind=closest
 #SBATCH --account=bche-delta-gpu
-#SBATCH -t 00:15:00
+#SBATCH -t 00:10:00
-#SBATCH --exclude=gpub029,gpub034,gpub038,gpub043,gpub074,gpub054,gpub012
+#SBATCH --exclude=gpub029,gpub034,gpub038,gpub043,gpub074,gpub054,gpub012,gpub044,gpub083
 #SBATCH --mail-user=asing271@illinois.edu
 #SBATCH --mail-type=BEGIN,END,FAIL
 
@@ -22,7 +22,7 @@ module load cuda
 echo -e "job $SLURM_JOBID is starting on `hostname`\n\n"
 
 # ── nsys system-level profile ─────────────────────────────────────
-srun nsys profile --stats=true ./test_gpt2 > test_gpt2_nsys_prof_req_0.out
+srun nsys profile --stats=true ./test_gpt2 > test_gpt2_nsys_prof_req_perf.out
 
 # ── ncu kernel-level profile ──────────────────────────────────────
 # ncu --set full -f -o analysis_file ./test_gpt2 > test_gpt2_ncu.out