diff --git a/src/kernels/matmul.cu b/src/kernels/matmul.cu index 2441dbe..ae27cc5 100644 --- a/src/kernels/matmul.cu +++ b/src/kernels/matmul.cu @@ -46,7 +46,7 @@ __global__ void Kernels::vec_scalar_sub( if (tid >= len) { return; } - d_out[tid] = d_src[tid] - d_scalar[0]; + d_out[tid] = d_src[tid] - *d_scalar; } __global__ void Kernels::vec_scalar_div( @@ -59,7 +59,7 @@ __global__ void Kernels::vec_scalar_div( if (tid >= len) { return; } - d_out[tid] = d_src[tid] / d_scalar[0]; + d_out[tid] = d_src[tid] / *d_scalar; } __global__ void Kernels::vec_exp( diff --git a/src/layers/activation.cu b/src/layers/activation.cu index 8171c8b..6d14e78 100644 --- a/src/layers/activation.cu +++ b/src/layers/activation.cu @@ -52,7 +52,7 @@ void Activation::activate(float* d_input) { // Subtract max value to improve numerical stability Kernels::vec_scalar_sub<<>>( - d_input, d_input, d_max, length + d_input, d_input, &d_max[0], length ); CUDA_CHECK(cudaGetLastError()); @@ -66,7 +66,7 @@ void Activation::activate(float* d_input) { Utils::sum(d_input, d_softmax_sum, length); Kernels::vec_scalar_div<<>>( - d_input, d_input, d_softmax_sum, length + d_input, d_input, &d_softmax_sum[0], length ); CUDA_CHECK(cudaGetLastError()); break;