From 0490dd77db7d3151ea3d719f52ba9ae416ec155b Mon Sep 17 00:00:00 2001 From: gabrielequatrana Date: Sun, 5 May 2024 17:39:44 +0200 Subject: [PATCH] ELG SpMM now working --- cuda/spgpu/kernels/ell_spmm_base.cuh | 7 +++---- cuda/spgpu/kernels/ell_spmm_base_template.cuh | 7 ++----- 2 files changed, 5 insertions(+), 9 deletions(-) diff --git a/cuda/spgpu/kernels/ell_spmm_base.cuh b/cuda/spgpu/kernels/ell_spmm_base.cuh index 9f76bfc0..da54794c 100644 --- a/cuda/spgpu/kernels/ell_spmm_base.cuh +++ b/cuda/spgpu/kernels/ell_spmm_base.cuh @@ -122,14 +122,13 @@ GEN_SPGPU_ELL_NAME(TYPE_SYMBOL) VALUE_TYPE beta, int baseIndex) { -// TODO VALUE_TYPE *px,*py,*pz; int cnt; int maxNForACall = max(handle->maxGridSizeX, THREAD_BLOCK*handle->maxGridSizeX); int maxShmemSz; maxShmemSz=getGPUSharedMemPerBlock(); - + //fprintf(stderr,"MaxSHmemSz %d \n",maxShmemSz); while (rows > maxNForACall) {//managing large vectors cnt = count; px = (VALUE_TYPE *) x; @@ -176,7 +175,7 @@ GEN_SPGPU_ELL_NAME(TYPE_SYMBOL) alpha, cM, rP, cMPitch, rPPitch, rS, rIdx, avgNnzPerRow, - maxNnzPerRow, maxNForACall, + maxNnzPerRow, rows, px, xPitch, beta, baseIndex); px += xPitch*MMBSZ; py += yPitch*MMBSZ; @@ -189,7 +188,7 @@ GEN_SPGPU_ELL_NAME(TYPE_SYMBOL) alpha, cM, rP, cMPitch, rPPitch, rS, rIdx, avgNnzPerRow, - maxNnzPerRow, maxNForACall, + maxNnzPerRow, rows, px, xPitch, beta, baseIndex); } diff --git a/cuda/spgpu/kernels/ell_spmm_base_template.cuh b/cuda/spgpu/kernels/ell_spmm_base_template.cuh index c5b39a95..60732acb 100644 --- a/cuda/spgpu/kernels/ell_spmm_base_template.cuh +++ b/cuda/spgpu/kernels/ell_spmm_base_template.cuh @@ -34,12 +34,9 @@ CONCAT(GEN_SPGPU_ELL_NAME(TYPE_SYMBOL), _krn) int i = threadIdx.x + blockIdx.x * (THREAD_BLOCK); if (i < rows) { - int j; + rS += i; rP += i; cM += i; - rP += rPPitch; - cM += cMPitch; - - int rowSize = rS[i]; + int rowSize = rS[0]; for (int k=0; k