Mathis/CUDANet
1
0
Fork 0
mirror of https://github.com/lordmathis/CUDANet.git synced 2025-11-06 01:34:22 +00:00
Code Issues Packages Projects Releases Wiki Activity

Change vector scalar math kernels

Browse Source
This commit is contained in:
Matúš Námešný
2024-04-28 19:02:57 +02:00
parent 69111f6cb1
commit 1f708f3ea3
2 changed files with 4 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
src/kernels/matmul.cu
View File

@@ -46,7 +46,7 @@ __global__ void Kernels::vec_scalar_sub(
if (tid >= len) { if (tid >= len) {
return; return;
} }
d_out[tid] = d_src[tid] - d_scalar[0]; d_out[tid] = d_src[tid] - *d_scalar;
} }
__global__ void Kernels::vec_scalar_div( __global__ void Kernels::vec_scalar_div(
@@ -59,7 +59,7 @@ __global__ void Kernels::vec_scalar_div(
if (tid >= len) { if (tid >= len) {
return; return;
} }
d_out[tid] = d_src[tid] / d_scalar[0]; d_out[tid] = d_src[tid] / *d_scalar;
} }
__global__ void Kernels::vec_exp( __global__ void Kernels::vec_exp(

4
src/layers/activation.cu
View File

@@ -52,7 +52,7 @@ void Activation::activate(float* d_input) {
// Subtract max value to improve numerical stability // Subtract max value to improve numerical stability
Kernels::vec_scalar_sub<<<gridSize, BLOCK_SIZE>>>( Kernels::vec_scalar_sub<<<gridSize, BLOCK_SIZE>>>(
d_input, d_input, d_max, length d_input, d_input, &d_max[0], length
); );
CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaGetLastError());
@@ -66,7 +66,7 @@ void Activation::activate(float* d_input) {
Utils::sum(d_input, d_softmax_sum, length); Utils::sum(d_input, d_softmax_sum, length);
Kernels::vec_scalar_div<<<gridSize, BLOCK_SIZE>>>( Kernels::vec_scalar_div<<<gridSize, BLOCK_SIZE>>>(
d_input, d_input, d_softmax_sum, length d_input, d_input, &d_softmax_sum[0], length
); );
CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaGetLastError());
break; break;