From 7e2e73a339137faefdc4ef6704e435c380728110 Mon Sep 17 00:00:00 2001 From: Aakarsh Kashyap Date: Fri, 1 May 2026 02:05:25 +0530 Subject: [PATCH 1/5] changed cudaDeviceSynchronize with cudaGetLastError and made memset async --- benchmarks/bench_softcuda.cpp | 2 +- src/backend_cpu/backprop/backprop_cuda_bridge.cu | 2 +- src/backend_gpu/math/add.cu | 2 +- src/backend_gpu/math/broadcast_add.cu | 2 +- src/backend_gpu/math/matmul.cu | 3 ++- src/backend_gpu/math/mean.cu | 2 +- src/backend_gpu/math/relu.cu | 2 +- src/backend_gpu/math/scalar_mul.cu | 2 +- src/backend_gpu/math/square.cu | 2 +- src/backend_gpu/math/sub.cu | 2 +- 10 files changed, 11 insertions(+), 10 deletions(-) diff --git a/benchmarks/bench_softcuda.cpp b/benchmarks/bench_softcuda.cpp index e9d27bb..5b2cf8e 100644 --- a/benchmarks/bench_softcuda.cpp +++ b/benchmarks/bench_softcuda.cpp @@ -99,7 +99,7 @@ static void bench_matmul() { header("Benchmark 2: Matmul 4096×4096"); const uint32_t M = 4096, K = 4096, N = 4096; - const uint32_t total_flops = 2 * M * K * N; + const size_t total_flops = 2 * M * K * N; float *A = new float[M * K]; float *B = new float[K * N]; for (uint32_t i = 0; i < M*K; i++) A[i] = (float)i * 0.001f; diff --git a/src/backend_cpu/backprop/backprop_cuda_bridge.cu b/src/backend_cpu/backprop/backprop_cuda_bridge.cu index 11032e4..1da5d2c 100644 --- a/src/backend_cpu/backprop/backprop_cuda_bridge.cu +++ b/src/backend_cpu/backprop/backprop_cuda_bridge.cu @@ -3,7 +3,7 @@ #include void soft_cuda_memset_zero(void *ptr, size_t bytes) { - cudaMemset(ptr, 0, bytes); + if (ptr) cudaMemsetAsync(ptr, 0, bytes); } void soft_cuda_memcpy_h2d(void *dst, const void *src, size_t bytes) { diff --git a/src/backend_gpu/math/add.cu b/src/backend_gpu/math/add.cu index 9b34c0a..b7cfa64 100644 --- a/src/backend_gpu/math/add.cu +++ b/src/backend_gpu/math/add.cu @@ -18,7 +18,7 @@ bool tensor_add_op_cuda(tensor_t *t,float *d_a, float *d_b, float *d_res) { int numBlocks = (t->nvalues + blockSize -1) / blockSize; add<<>>(d_a,d_b,d_res,t->nvalues); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA Add Kernel Failed: %s\n", cudaGetErrorString(err)); diff --git a/src/backend_gpu/math/broadcast_add.cu b/src/backend_gpu/math/broadcast_add.cu index c34efa7..952315f 100644 --- a/src/backend_gpu/math/broadcast_add.cu +++ b/src/backend_gpu/math/broadcast_add.cu @@ -22,7 +22,7 @@ bool tensor_broadcast_add_op_cuda(tensor_t *t, float *d_a, float *d_b, int block = 256; int grid = ((int)total + block - 1) / block; broadcast_add_kernel<<>>(d_a, d_b, d_res, rows, cols); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA BroadcastAdd Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/matmul.cu b/src/backend_gpu/math/matmul.cu index 9c02673..7238c98 100644 --- a/src/backend_gpu/math/matmul.cu +++ b/src/backend_gpu/math/matmul.cu @@ -65,8 +65,9 @@ bool tensor_mul_op_cuda(tensor_t *t, float *d_a, float *d_b, float *d_res) { debug("tensor_mul_op_cuda: cublasSgemm failed (%d)\n", (int)stat); return false; } + cudaError_t err = cudaGetLastError(); + - cudaError_t err = cudaDeviceSynchronize(); if (err != cudaSuccess) { debug("CUDA Matmul Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/mean.cu b/src/backend_gpu/math/mean.cu index a40afe9..0ed5ace 100644 --- a/src/backend_gpu/math/mean.cu +++ b/src/backend_gpu/math/mean.cu @@ -67,7 +67,7 @@ bool tensor_mean_op_cuda(tensor_t *t, float *d_a, float *d_res) { cudaFree(d_partial); cudaFree(d_partial2); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA Mean Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/relu.cu b/src/backend_gpu/math/relu.cu index cb440be..06b0c8f 100644 --- a/src/backend_gpu/math/relu.cu +++ b/src/backend_gpu/math/relu.cu @@ -13,7 +13,7 @@ bool tensor_relu_op_cuda(tensor_t *t, float *d_a, float *d_res) { int block = 256; int grid = ((int)t->nvalues + block - 1) / block; relu_kernel<<>>(d_a, d_res, t->nvalues); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA ReLU Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/scalar_mul.cu b/src/backend_gpu/math/scalar_mul.cu index 0995107..8438a44 100644 --- a/src/backend_gpu/math/scalar_mul.cu +++ b/src/backend_gpu/math/scalar_mul.cu @@ -17,7 +17,7 @@ bool tensor_scalar_mul_op_cuda(tensor_t *t, float *d_a, int block = 256; int grid = ((int)t->nvalues + block - 1) / block; scalar_mul_kernel<<>>(d_a, d_res, s, t->nvalues); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA ScalarMul Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/square.cu b/src/backend_gpu/math/square.cu index 4442e37..3dce92c 100644 --- a/src/backend_gpu/math/square.cu +++ b/src/backend_gpu/math/square.cu @@ -13,7 +13,7 @@ bool tensor_square_op_cuda(tensor_t *t, float *d_a, float *d_res) { int block = 256; int grid = ((int)t->nvalues + block - 1) / block; square_kernel<<>>(d_a, d_res, t->nvalues); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA Square Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/sub.cu b/src/backend_gpu/math/sub.cu index 0913912..b4def88 100644 --- a/src/backend_gpu/math/sub.cu +++ b/src/backend_gpu/math/sub.cu @@ -13,7 +13,7 @@ bool tensor_sub_op_cuda(tensor_t *t, float *d_a, float *d_b, float *d_res) { int block = 256; int grid = ((int)t->nvalues + block - 1) / block; sub_kernel<<>>(d_a, d_b, d_res, t->nvalues); - cudaError_t err = cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { debug("CUDA Sub Kernel Failed: %s\n", cudaGetErrorString(err)); return false; From 818f035839f877be1b9a8f695299b280275ea58f Mon Sep 17 00:00:00 2001 From: Aakarsh Kashyap Date: Fri, 1 May 2026 02:24:35 +0530 Subject: [PATCH 2/5] Some sync issue were still left removed them i guess testing suite will be written next i guess --- src/backend_gpu/backprop/backprop_gpu.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/backend_gpu/backprop/backprop_gpu.cu b/src/backend_gpu/backprop/backprop_gpu.cu index b4e3782..f348f0a 100644 --- a/src/backend_gpu/backprop/backprop_gpu.cu +++ b/src/backend_gpu/backprop/backprop_gpu.cu @@ -306,7 +306,7 @@ bool backprop_gpu_dispatch(execution_node_t *node, return false; } - err = cudaDeviceSynchronize(); + err = cudaGetLastError(); if (err != cudaSuccess) { debug("backprop_gpu_dispatch: CUDA error: %s\n", cudaGetErrorString(err)); return false; @@ -325,6 +325,6 @@ extern "C" bool tensor_sgd_gpu(float *d_w, float *d_g, float lr, uint32_t n) { int block = 256; int grid = ((int)n + block - 1) / block; gpu_sgd_k<<>>(d_w, d_g, lr, n); - return cudaDeviceSynchronize() == cudaSuccess; + return cudaGetLastError() == cudaSuccess; } From a40d9381f2fbf35eb7ffa0f5b0ebb60a1960d0a7 Mon Sep 17 00:00:00 2001 From: Aakarsh Kashyap Date: Fri, 1 May 2026 02:47:47 +0530 Subject: [PATCH 3/5] Added definition gaurd against cudasync and cudalasterror for debug build --- src/backend_gpu/backprop/backprop_gpu.cu | 13 ++++++++++++- src/backend_gpu/math/add.cu | 8 +++++++- src/backend_gpu/math/broadcast_add.cu | 11 ++++++++++- src/backend_gpu/math/matmul.cu | 8 +++++++- src/backend_gpu/math/mean.cu | 8 +++++++- src/backend_gpu/math/relu.cu | 8 +++++++- src/backend_gpu/math/scalar_mul.cu | 8 +++++++- src/backend_gpu/math/square.cu | 8 +++++++- src/backend_gpu/math/sub.cu | 8 +++++++- 9 files changed, 71 insertions(+), 9 deletions(-) diff --git a/src/backend_gpu/backprop/backprop_gpu.cu b/src/backend_gpu/backprop/backprop_gpu.cu index f348f0a..35478e9 100644 --- a/src/backend_gpu/backprop/backprop_gpu.cu +++ b/src/backend_gpu/backprop/backprop_gpu.cu @@ -306,7 +306,11 @@ bool backprop_gpu_dispatch(execution_node_t *node, return false; } +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("backprop_gpu_dispatch: CUDA error: %s\n", cudaGetErrorString(err)); return false; @@ -325,6 +329,13 @@ extern "C" bool tensor_sgd_gpu(float *d_w, float *d_g, float lr, uint32_t n) { int block = 256; int grid = ((int)n + block - 1) / block; gpu_sgd_k<<>>(d_w, d_g, lr, n); - return cudaGetLastError() == cudaSuccess; +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif + return err == cudaSuccess; } diff --git a/src/backend_gpu/math/add.cu b/src/backend_gpu/math/add.cu index b7cfa64..00a24f3 100644 --- a/src/backend_gpu/math/add.cu +++ b/src/backend_gpu/math/add.cu @@ -18,7 +18,13 @@ bool tensor_add_op_cuda(tensor_t *t,float *d_a, float *d_b, float *d_res) { int numBlocks = (t->nvalues + blockSize -1) / blockSize; add<<>>(d_a,d_b,d_res,t->nvalues); - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("CUDA Add Kernel Failed: %s\n", cudaGetErrorString(err)); diff --git a/src/backend_gpu/math/broadcast_add.cu b/src/backend_gpu/math/broadcast_add.cu index 952315f..eb15309 100644 --- a/src/backend_gpu/math/broadcast_add.cu +++ b/src/backend_gpu/math/broadcast_add.cu @@ -22,7 +22,16 @@ bool tensor_broadcast_add_op_cuda(tensor_t *t, float *d_a, float *d_b, int block = 256; int grid = ((int)total + block - 1) / block; broadcast_add_kernel<<>>(d_a, d_b, d_res, rows, cols); - cudaError_t err = cudaGetLastError(); + +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif + + if (err != cudaSuccess) { debug("CUDA BroadcastAdd Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/matmul.cu b/src/backend_gpu/math/matmul.cu index 7238c98..d258a85 100644 --- a/src/backend_gpu/math/matmul.cu +++ b/src/backend_gpu/math/matmul.cu @@ -65,7 +65,13 @@ bool tensor_mul_op_cuda(tensor_t *t, float *d_a, float *d_b, float *d_res) { debug("tensor_mul_op_cuda: cublasSgemm failed (%d)\n", (int)stat); return false; } - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { diff --git a/src/backend_gpu/math/mean.cu b/src/backend_gpu/math/mean.cu index 0ed5ace..c8342d4 100644 --- a/src/backend_gpu/math/mean.cu +++ b/src/backend_gpu/math/mean.cu @@ -67,7 +67,13 @@ bool tensor_mean_op_cuda(tensor_t *t, float *d_a, float *d_res) { cudaFree(d_partial); cudaFree(d_partial2); - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("CUDA Mean Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/relu.cu b/src/backend_gpu/math/relu.cu index 06b0c8f..41a0c5d 100644 --- a/src/backend_gpu/math/relu.cu +++ b/src/backend_gpu/math/relu.cu @@ -13,7 +13,13 @@ bool tensor_relu_op_cuda(tensor_t *t, float *d_a, float *d_res) { int block = 256; int grid = ((int)t->nvalues + block - 1) / block; relu_kernel<<>>(d_a, d_res, t->nvalues); - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("CUDA ReLU Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/scalar_mul.cu b/src/backend_gpu/math/scalar_mul.cu index 8438a44..580b0de 100644 --- a/src/backend_gpu/math/scalar_mul.cu +++ b/src/backend_gpu/math/scalar_mul.cu @@ -17,7 +17,13 @@ bool tensor_scalar_mul_op_cuda(tensor_t *t, float *d_a, int block = 256; int grid = ((int)t->nvalues + block - 1) / block; scalar_mul_kernel<<>>(d_a, d_res, s, t->nvalues); - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("CUDA ScalarMul Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/square.cu b/src/backend_gpu/math/square.cu index 3dce92c..e67a283 100644 --- a/src/backend_gpu/math/square.cu +++ b/src/backend_gpu/math/square.cu @@ -13,7 +13,13 @@ bool tensor_square_op_cuda(tensor_t *t, float *d_a, float *d_res) { int block = 256; int grid = ((int)t->nvalues + block - 1) / block; square_kernel<<>>(d_a, d_res, t->nvalues); - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("CUDA Square Kernel Failed: %s\n", cudaGetErrorString(err)); return false; diff --git a/src/backend_gpu/math/sub.cu b/src/backend_gpu/math/sub.cu index b4def88..bb24546 100644 --- a/src/backend_gpu/math/sub.cu +++ b/src/backend_gpu/math/sub.cu @@ -13,7 +13,13 @@ bool tensor_sub_op_cuda(tensor_t *t, float *d_a, float *d_b, float *d_res) { int block = 256; int grid = ((int)t->nvalues + block - 1) / block; sub_kernel<<>>(d_a, d_b, d_res, t->nvalues); - cudaError_t err = cudaGetLastError(); +cudaError_t err = cudaSuccess; + +#ifdef SC_DEBUG + err = cudaDeviceSynchronize(); +#else + err = cudaGetLastError(); +#endif if (err != cudaSuccess) { debug("CUDA Sub Kernel Failed: %s\n", cudaGetErrorString(err)); return false; From 0dd00c8253539a49a4849f3bc1677ce55e0cff71 Mon Sep 17 00:00:00 2001 From: Aakarsh Kashyap Date: Fri, 1 May 2026 03:12:53 +0530 Subject: [PATCH 4/5] Update src/backend_cpu/backprop/backprop_cuda_bridge.cu Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- src/backend_cpu/backprop/backprop_cuda_bridge.cu | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/src/backend_cpu/backprop/backprop_cuda_bridge.cu b/src/backend_cpu/backprop/backprop_cuda_bridge.cu index 1da5d2c..b7ab360 100644 --- a/src/backend_cpu/backprop/backprop_cuda_bridge.cu +++ b/src/backend_cpu/backprop/backprop_cuda_bridge.cu @@ -3,7 +3,11 @@ #include void soft_cuda_memset_zero(void *ptr, size_t bytes) { - if (ptr) cudaMemsetAsync(ptr, 0, bytes); + if (ptr) { + cudaError_t err = cudaMemsetAsync(ptr, 0, bytes); + if (err != cudaSuccess) + debug("cudaMemsetAsync failed: %s\n", cudaGetErrorString(err)); + } } void soft_cuda_memcpy_h2d(void *dst, const void *src, size_t bytes) { From 8620e3b3c3fbaa65bcbb3696cb1128fced850968 Mon Sep 17 00:00:00 2001 From: Aakarsh Kashyap Date: Fri, 1 May 2026 03:17:15 +0530 Subject: [PATCH 5/5] Update benchmarks/bench_softcuda.cpp Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> --- benchmarks/bench_softcuda.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/bench_softcuda.cpp b/benchmarks/bench_softcuda.cpp index 5b2cf8e..a924e72 100644 --- a/benchmarks/bench_softcuda.cpp +++ b/benchmarks/bench_softcuda.cpp @@ -99,7 +99,7 @@ static void bench_matmul() { header("Benchmark 2: Matmul 4096×4096"); const uint32_t M = 4096, K = 4096, N = 4096; - const size_t total_flops = 2 * M * K * N; + const size_t total_flops = (size_t)2 * M * K * N; float *A = new float[M * K]; float *B = new float[K * N]; for (uint32_t i = 0; i < M*K; i++) A[i] = (float)i * 0.001f;