Skip to content

Improve CUDA synchronization and make memset asynchronous #40

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
souls-syntax merged 5 commits into from
Apr 30, 2026
Merged

Improve CUDA synchronization and make memset asynchronous #40

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
7e2e73a
changed cudaDeviceSynchronize with cudaGetLastError and made memset a…
souls-syntax Apr 30, 2026
818f035
Some sync issue were still left removed them i guess testing suite wi…
souls-syntax Apr 30, 2026
a40d938
Added definition gaurd against cudasync and cudalasterror for debug b…
souls-syntax Apr 30, 2026
0dd00c8
Update src/backend_cpu/backprop/backprop_cuda_bridge.cu
souls-syntax Apr 30, 2026
8620e3b
Update benchmarks/bench_softcuda.cpp
souls-syntax Apr 30, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion benchmarks/bench_softcuda.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 = (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;
Expand Down
6 changes: 5 additions & 1 deletion src/backend_cpu/backprop/backprop_cuda_bridge.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,11 @@
#include <cassert>

void soft_cuda_memset_zero(void *ptr, size_t bytes) {
cudaMemset(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) {
Expand Down
13 changes: 12 additions & 1 deletion src/backend_gpu/backprop/backprop_gpu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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<<<grid, block>>>(d_w, d_g, lr, n);
return cudaDeviceSynchronize() == cudaSuccess;
cudaError_t err = cudaSuccess;

#ifdef SC_DEBUG
err = cudaDeviceSynchronize();
#else
err = cudaGetLastError();
#endif
return err == cudaSuccess;
}

8 changes: 7 additions & 1 deletion src/backend_gpu/math/add.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<numBlocks, blockSize>>>(d_a,d_b,d_res,t->nvalues);

cudaError_t err = cudaDeviceSynchronize();
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));
Expand Down
11 changes: 10 additions & 1 deletion src/backend_gpu/math/broadcast_add.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<grid, block>>>(d_a, d_b, d_res, rows, cols);
cudaError_t err = cudaDeviceSynchronize();

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;
Expand Down
9 changes: 8 additions & 1 deletion src/backend_gpu/math/matmul.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,8 +65,15 @@ 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 = cudaSuccess;

#ifdef SC_DEBUG
err = cudaDeviceSynchronize();
#else
err = cudaGetLastError();
#endif


cudaError_t err = cudaDeviceSynchronize();
if (err != cudaSuccess) {
debug("CUDA Matmul Kernel Failed: %s\n", cudaGetErrorString(err));
return false;
Expand Down
8 changes: 7 additions & 1 deletion src/backend_gpu/math/mean.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 = cudaDeviceSynchronize();
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;
Expand Down
8 changes: 7 additions & 1 deletion src/backend_gpu/math/relu.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<grid, block>>>(d_a, d_res, t->nvalues);
cudaError_t err = cudaDeviceSynchronize();
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;
Expand Down
8 changes: 7 additions & 1 deletion src/backend_gpu/math/scalar_mul.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<grid, block>>>(d_a, d_res, s, t->nvalues);
cudaError_t err = cudaDeviceSynchronize();
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;
Expand Down
8 changes: 7 additions & 1 deletion src/backend_gpu/math/square.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<grid, block>>>(d_a, d_res, t->nvalues);
cudaError_t err = cudaDeviceSynchronize();
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;
Expand Down
8 changes: 7 additions & 1 deletion src/backend_gpu/math/sub.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<<<grid, block>>>(d_a, d_b, d_res, t->nvalues);
cudaError_t err = cudaDeviceSynchronize();
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;
Expand Down
Loading