Minor fixed

Author: RobBa

readme.md

@@ -22,6 +22,7 @@ For some examples on Python interface, see tests/python.
 - **Python Interface**: Seamless integration via Boost.Python
 - **Clean Architecture**: Modular design, maintainable and extensible
 - **CI/CD**: Automated testing with GTest and GitHub Actions
+- **CUDA implementation**: Full end-to-end CUDA implementation for better HW utilization (optimization in progress)
 
 ## Tech Stack
 

src/backend/module/activation_functions/cuda/activations.cu

@@ -330,8 +330,7 @@ namespace cuda_impl {
 
     // TODO: use some static struct here to prevent this guy from keeping on re-allocating memory
     ftype* maxValues;
-    const tensorSize_t nMaxValues = nStrides;
-    cudaErrchk(cudaMalloc(&maxValues, nMaxValues * sizeof(ftype)));
+    cudaErrchk(cudaMalloc(&maxValues, nStrides * sizeof(ftype)));
 
     static const auto warpSizeT2 = 2 * DeviceProperties::getWarpSize(); // TODO: can this be a problem in a multi-GPU setting?
     if(stride <= warpSizeT2) {
@@ -340,7 +339,7 @@ namespace cuda_impl {
       const int threadsPerBlock = 256;
       const int blocks = (in.getSize() + threadsPerBlock - 1) / threadsPerBlock;
 
-      if(stride == 2) {
+      if(stride <= 2) {
         findMaxKernelOneWarp<1> <<<blocks, threadsPerBlock, threadsPerBlock * sizeof(ftype)>>>(maxValues, in.getData(), stride, in.getSize());
       }
       else if(stride <= 4) {

src/backend/shared/cuda/common_softmax.cuh

@@ -68,15 +68,17 @@ namespace cuda_impl {
   static __global__ void findMaxKernelOneWarp(ftype* const res, const ftype* const input, const tensorSize_t stride, const tensorSize_t size) {
     assert(blockDim.x % 32 == 0);
 
-    int gid = blockIdx.x * blockDim.x + threadIdx.x;
-    if(gid >= size)
-      return;
-
-    int tid = threadIdx.x;
+    const int tid = threadIdx.x;
+    const int gid = blockIdx.x * blockDim.x + threadIdx.x;
+    const bool isNotPadded = gid < size;
     extern __shared__ ftype smem[];
-    smem[tid] = input[gid];
+    smem[tid] = isNotPadded ? input[gid] : -INFINITY;
     __syncthreads();
 
+    if(!isNotPadded) {
+      return;
+    }
+
     volatile ftype* const start = smem + (tid / stride) * stride;
     const int offset = gid % stride;
     warpMaxReduce<maxoffset>(start, stride, offset);