diff --git a/benchmarks/bench_softcuda.cpp b/benchmarks/bench_softcuda.cpp index e9d27bb..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 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; diff --git a/src/backend_cpu/backprop/backprop_cuda_bridge.cu b/src/backend_cpu/backprop/backprop_cuda_bridge.cu index 11032e4..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) { - 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) { diff --git a/src/backend_gpu/backprop/backprop_gpu.cu b/src/backend_gpu/backprop/backprop_gpu.cu index b4e3782..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 cudaDeviceSynchronize() == 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 9b34c0a..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 = 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)); diff --git a/src/backend_gpu/math/broadcast_add.cu b/src/backend_gpu/math/broadcast_add.cu index c34efa7..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 = 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; diff --git a/src/backend_gpu/math/matmul.cu b/src/backend_gpu/math/matmul.cu index 9c02673..d258a85 100644 --- a/src/backend_gpu/math/matmul.cu +++ b/src/backend_gpu/math/matmul.cu @@ -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; diff --git a/src/backend_gpu/math/mean.cu b/src/backend_gpu/math/mean.cu index a40afe9..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 = 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; diff --git a/src/backend_gpu/math/relu.cu b/src/backend_gpu/math/relu.cu index cb440be..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 = 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; diff --git a/src/backend_gpu/math/scalar_mul.cu b/src/backend_gpu/math/scalar_mul.cu index 0995107..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 = 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; diff --git a/src/backend_gpu/math/square.cu b/src/backend_gpu/math/square.cu index 4442e37..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 = 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; diff --git a/src/backend_gpu/math/sub.cu b/src/backend_gpu/math/sub.cu index 0913912..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 = 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;