From 545063223720505e4da4ef1bdea176d8006bf4f6 Mon Sep 17 00:00:00 2001
From: Christopher <christopher.paul.dryden@gmail.com>
Date: Thu, 30 May 2024 03:01:52 +0000
Subject: [PATCH] Removed unnecesary shared memory due to blockreduce using
 static defined shared memory

---
 dev/cuda/classifier_fused.cu | 2 +-
 train_gpt2.cu                | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/dev/cuda/classifier_fused.cu b/dev/cuda/classifier_fused.cu
index 2125b87..5b6c986 100644
--- a/dev/cuda/classifier_fused.cu
+++ b/dev/cuda/classifier_fused.cu
@@ -664,7 +664,7 @@ void fused_classifier5(float* dlogits, float* losses,
                       int B, int T, int V, int P, int block_size) {
     const int N = B * T;
     const int grid_size = N;
-    fused_classifier_kernel5<true,false><<<grid_size, block_size, 512>>>((floatX*)dlogits, (floatX*)losses, NULL, (floatX*)logits, (floatX*)dlosses, targets, B, T, V, P);
+    fused_classifier_kernel5<true,false><<<grid_size, block_size>>>((floatX*)dlogits, (floatX*)losses, NULL, (floatX*)logits, (floatX*)dlosses, targets, B, T, V, P);
     cudaCheck(cudaGetLastError());
 }
 
diff --git a/train_gpt2.cu b/train_gpt2.cu
index 5918e21..053c2e3 100644
--- a/train_gpt2.cu
+++ b/train_gpt2.cu
@@ -1899,7 +1899,7 @@ void fused_classifier(Type* logits, Type* losses,
     const int block_size = 1024;
     const int N = B * T;
     const int grid_size = N;
-    fused_classifier_kernel5<<<grid_size, block_size, 512>>>(logits, losses, (floatX*)NULL, dloss, targets, B, T, V, P);
+    fused_classifier_kernel5<<<grid_size, block_size>>>(logits, losses, (floatX*)NULL, dloss, targets, B, T, V, P);
     cudaCheck(cudaGetLastError());
 }