From 1f708f3ea395387cb5137d3c0fde0819a30e5a1b Mon Sep 17 00:00:00 2001
From: LordMathis <matus@namesny.com>
Date: Sun, 28 Apr 2024 19:02:57 +0200
Subject: [PATCH] Change vector scalar math kernels

---
 src/kernels/matmul.cu    | 4 ++--
 src/layers/activation.cu | 4 ++--
 2 files changed, 4 insertions(+), 4 deletions(-)

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<<<gridSize, BLOCK_SIZE>>>(
-                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<<<gridSize, BLOCK_SIZE>>>(
-                d_input, d_input, d_softmax_sum, length
+                d_input, d_input, &d_softmax_sum[0], length
             );
             CUDA_CHECK(cudaGetLastError());
             break;