SAS: Structured Activation Sparsification
Authors: Yusuke Sekikawa
Shingo Yashima
Venue: ICLR 2024 poster (2024)
Keywords: Structured
Sparse
Pruning
Projection
Primary area: general machine learning (i.e., none of the above)
TL;DR: Utilize structured sparsity in activation by sparse projection aiming to realized efficient vectorized matmul

Abstract: Wide networks usually yield better accuracy than their narrower counterpart at the expense of the massive $\texttt{mult}$ cost.
To break this tradeoff, we advocate a novel concept of $\textit{Structured Activation Sparsification}$, dubbed SAS, which boosts accuracy without increasing computation by utilizing the projected sparsity in activation maps with a specific structure. 
Concretely, the projected sparse activation is allowed to have N nonzero value among M consecutive activations.
Owing to the local structure in sparsity, the wide $\texttt{matmul}$ between a dense weight and the sparse activation is executed as an equivalent narrow $\texttt{matmul}$ between a dense weight and dense activation, which is compatible with NVIDIA's $\textit{SparseTensorCore}$ developed for the N:M structured sparse weight.
In extensive experiments, we demonstrate that increasing sparsity monotonically improves accuracy (up to 7% on CIFAR10) without increasing the $\texttt{mult}$ count.
Furthermore, we show that structured sparsification of $\textit{activation}$ scales better than that of $\textit{weight}$ given the same computational budget.

Published as a conference paper at ICLR 2024
SAS: STRUCTURED ACTIVATION SPARSIFICATION
Yusuke Sekikwa, Shingo Yashima DENSO IT Lab. Inc., Tokyo, Japan Sasueeseeiikaee yashima... waw.aooooooor.diillbb...b......a
ABSTRACT
Wide networks usually yield better accuracy than their narrower counter- part at the expense of the massive mult cost. To break this tradeoff, we ad- vocate a novel concept of Structured Activation Spersification, dubbed SAS, which boosts accuracy without increasing computation by utilizing the pro- jected sparsity in activation maps with a specific structure. Concretely, the projected sparse activation is allowed to have N nonzero value among M consecutive aclivalions. Owing to the local structure sparsily, the wide matmul between dense weight and the sparse activation is executed an equivalent narrow matmul between a dense weight and dense activation, which is compatible with NVIDIA's Sparse Tensor Core developed for the : .M structured sparse weight. In extensive experiments, we demonstrate that increasing sparsity monotonically improves accuracy (up to 7% on CIFARIO) without increasing the mult count. Furthermore, we show that structured sparsification of aclivation scales better than that of weight given the same computational budget. hhtps: //t::////ttbb
INTRODUCTION
Modern deep neural networks (DNNs) consist of numerous matrix multiplications between input activation and weight matrices, making it challenging for DNNs to deploy resource- constrained edge devices due to the enormous mult count. One could reduce the computa- tion by narrowing the width of the activation channel; however, trades off accuracy.
Researchers explore the way to speed up DNN inference from various perspectives such as quantization (Esser al.) 2019), knowledge distillation (Hinton efficient model architectures Liu et al.] 2022; Howard et al. 2017). and structured weight
2015 designing
sparsification (Wang et 2019 Tan al. 2020: Li al 2022 2016 Wen et al. 2016). This
paper explores the unexplored spectrum, Structured Activation Sparsifieation, positioned at: the opposite end of the structured weight sparsífication.
WEIGHT SPARSIFICATION
One could obtain an efficient model by removing unimportant weights from wide, high- capacity model. It can be categorized into unstructured and structured.
Unstructured weight sparsification prunes individual weights independently based on the criteria for the importance of the weight connection (c.g., by their magnitude) (Frankle 2019 Wortsman et al.) 2019). Although could achieve an extremely sparse network without sacrilicing accuracy, the sparsity does not translate to the practical speedup on poctor-processee
& Carbin 2018 Yang
architectures (e-g., GPU) because it has no structure the sparsity pattern, making the parallel execution difficult (NVIDIA 2020). In fact, even when a matrix is highly sparse (e.g., >95% zero), wall-clock times for the sparse matmul on GPU (CuSPARSE) is still slower than dense operation (cuBLAS) (Shi ct al. 2020
Coarse- -raiindd structured weight sparsification removes unimportant weights such that the sparsified operation can be vectorized, which includes block pruning (Wang al. 2019 Chen et al. 2022), kernel shape sparsity (Tan et al. 2020), and channel/ /fitte pruning
Li et. al.] 2022
2016 Wen et al. 2016). The resultant network's theoretical mult count.

---

Published as a conference paper at ICLR 2024
directly translates to actual speedup on commodity vector-type hardware; however, it hurts model performance more severely than the unstructured ine-grained sparsity.
Fine-grained structured rremi-structuredd weight sparsification aims at the best of both worlds. The N M structured weights sparsification (SWS) utilized the local sparse structure in weight to realize the efficient parallel exccution of sparse matmul. The struc- ture indicates that the sparse weight has N nonzero value for every continuous M clement (section 5). NVIDIA commercialized this idea in Sparse Tensor Core It incorporates this local structure by the vectorized local switch, which selects corresponding N elements from M consecutive activation, realizing the dot-product using only N/M mult count. They demonstrate the actual speedup of SWS close to the theoretical gain. This real speedup on the wide-spread hardware using local sparsity patterns motivates us to explore the oppo- site: structured sparsity in activation, expecting it to deliver better accuracy/ speed tradeoff. Hereafter, we call the N:M structured pattern as structured for brevity.
ACTIVATION SPARSIFICATION
Sparsity also exists in activation; however, unlike sparsity in weight, much less attention has been given to exploiting sparsity in activation for DNN speedup.
Unstructured activation sparsification has been utilized to speed up the DNN inference.
In many settings, DNNs' activation tends to be naturally sparse by the activation functions, c.g., rectified linear unit (ReLU). Most of the prior works utilized the activation sparsity by
their novel hardware architecture (Zhu et al.) 2022 Albericio et al. 2016 Rhu et al., 2018
Han el. al., 2016 Parashar et. al. 2017 Park et al. 2017 Georgiadis 2019 Wang al.
2021 (Kurtz el al. 2020) demonstrate the wall-clock speedup on a general-purpent CPU.
However, the sparsity induced by activation functions such ReLU depends its input; therefore, the resultant sparsity pattern is unstructured. This unstructured nature makes it difficult to parallelize the operation on vector-type processors such as GPU (Shi al.) 2020 This incompatibility with the popular vector-type processors partly explains why activation sparsity has been less explored than weight sparsity.
Structured Activation Sparsification (SAS) is an unexplored area that we expect to pioneer the new frontier for accuracy sauuurccy/ppeee speed tradeoff in rector-processen architecture. Our contribution is two-fold: WW introduce the idea of SAS, open up new research area, and We invent mechanism called structured sparse projection to construct SAS, realized vectorized matmul with sparse activation on GPU for the first time. Realizing structured sparsity, i.e., controlling the nonzero counts within consecutive activation elements, is diffi- cult because the activation value depends on its input. We realized the structured sparsity in activation by implicitly projecting the input narrow/dens activation into higher dimen- sional space such that the projected activation has a local sparsity pattern. This contrasts to pruning as in SWS, which sparsify the value (weight) to be fixed after training. SAS yields better accuracy by utilizing the increased flexibility of the wide weight while keeping the same mult count. In other words, SAS selects the oppropriate weight depending on its input using the activation's sparsity pattern. Importantly, owing the local structure of the projected sparsity, SAS realized the vectorized matmul with sparse activation on typical GPU by utilizing the local routing mechanism originally developed for SWS.
Our experiments (section 3 demonstrate that increased sparsity in the SAS network en- hances its capacity and accuracy monotonically without increasing mult count; furthermore, it performs much better than the SWS for the same speed.
2 STRUCTURED ACTIVATION SPARSIFICATION
A matrix multiplication (matmul) is a fundamental building block of neural networks. This study mainly focuses on the matmul that appears in convolution ? Convolution with the activation I CC and weight W E &.CxCo is equivalent to matmul between
Ampere and beyond support this for 2:4 and 1:2 sparsity (NVIDIA 2020).
"The same discussion be applied the matmul in most DNN elements, such attention.
2

---

Published as a conference paper at ICLR 2024
Input activation
Structured sparse activation
Impilatt Projection
Input Activation
octivation Indices
111
Compress
Dot Product
Product
Activation
Weight
Output activation
Dense Muliplication
Mullipicaaion (1:4)
Figure 1: SAS overview. Conventional matmul between a dense/narrow weight and a dense/- narrow activation has lower expressive power (left). The matmul between dense/wit wide weight and structurally sparse/ /wide activation composed by implicit sparse projection (SAS) has higher pressive power while requiring the same mult count and comparable inference time (right).
unfolded activation X Ci and the reshaped weight W R.CGARR
(1) For dense weight W and dense activation X, the matmul consumes C,kk x C, x HW O. x C, x HW mult count. One can reduce the computation by narrowing the network width; mult count is reduced by factor a by narrowing the number of input and output channels by a. Unfortunately, but apparently, this trades off accuracy.
0 = + W W..
2.1 STRUCTURED SPARSITY IN ACTIVATION
We study an unxxpored area, exploiting structured sparsity in activation to realize a neural network model that is not only low FLOPS but also vectorizable. Concretely, the sparse activation is allowed to have N nonzero value among M consecutive activations. Owing to the sparsity, a matmul between sparse/ /wide activation X and dense/ wide weight W could be reduced into the matmul between dense/ narrow activation X and dense/narrow weight W as: O W = WX. More importantly, it realizes the parallel vectorized execution owing to the local sparsity pattern, specifically W is constructed by locally routing the weight elements from the sparse weight W that correspond to the nonzero activation elements in X. Actually, this operation is compatible with the commonly used GPU (section 2.3
2.2 SAS BY SPARSE PROJECTION
How can we construct such structurally sparse activation? Realizing structured sparsity (controlling the number of nonzero counts within consecutive activation elements) is chal- lenging because, unlike the case of SWS, the activation changes depending on the input. Therefore, existing methods utilizing activation sparsity unstruchured, such as induced by ReLU (Shi et al. 2020), which is hard to vectorize on GPU. One could construct a struc- turally sparse activation by maz-pooling along the blocked channel dimension; however, we can not expect this to achieve a good accuracy /sppeed tradeoff. Because @we can not control the input-dependant activation to have one (N) significant. value among M consec- utive element (otherwise, it removes crucial information); @ii discards most of the value in activation map (e.g., 94% when M-16) that is computed using valuable resources.
To realize the structured sparsity, we propose structured sparse projection S as follows:
S X X Rxx X E R M/NG
In the SAS projection S, single element in the source activation X is projected into one of an M consecutive elements in target sparse activation X indexed by I. In this study, use the sign of log,M noighbor activation to compute the index I,, for j-th input channel and w-th spacial location as follows:
Jog, M-1
Ij =
Σ
(Xftiau > 2)
(3)
i=0
3

---

Published as a conference paper at ICLR 2024
Input astin
Bruciured spanse activsiinn
sparse matadd (ochvator)
Sparse Projetton
HOOOOR activation Veriens
Posto matcting coement
MC.
Combress
Xpt
Poute
Der Product
Accumuttor
pup activation
Structured sparee
Weight
Figure 2: SAS mechanism. The matmul between sparse/ /wide activation X and dense wide weight W could be veclorized on Sparse Tensor Core by selecting elements from the weight W that correspond to the nonzero values in the activation X, skipping the unnecessary mult by zero (left). A ddense/narron activation X RCxxHW is projected to structurally sparse/ wide activation X R MeMMxxHWW (here M=4). Local index for projecting the activation is computed by using the sign of the log,M consecutive activation using eq. 3 (right). Note; we do not compute the sparse activation X explicitly during inference; instead, and I is directly passed to the processor.
1. def matmul(W. kSpars==2); # Activation X and Weight W. kSpame=2 indicates 50% sparsity beta, gamma, alpha W.sSap, W.haae[[]] X.shape[1]
3 Y 2etos[beta, alpba), dtype= Roat32") Outpus
4 E zerosi zeros([gammal / //ssp//s////22/ alpha), dtype='uintt")) | 0xd
Compute Index by evaluating agu bittbit (eg. (4)) and Convert 4-bit Index
[slgnbit(X)] Use
7 E bitwise or(E [0::.:). E (1:2))+<4))
8 E. reodered LEke array(mumpy load () aopp/foomaa((, k))) # need re-oder (supp. E, fg. A2)
Copy compresed activation to SparseTensorOCe compressed aevos acros( (alphaagaaaaaaa auda.runtime uutimme.mmmmpp... memepy(X compressed data.ptr. X. data.ptr, (lutXalphaggamma/ ganmma//SSparre)-4 ...)
(1Ppar)))(d+1111) dtype- "uints") # value followed by index
11
12
cuda.runtime. memepy(X compressed data dasta.prrdast size, E reodered data, (alpha-gamma/ /kksparree//
Execute matmul: W & X by calling cusparseLtbetaatun compressed data.ptr. ... Y.data.ptr....
cusparselt.muie mat.rell.mmmmml..... .... W.data.prr,
15
resurn
General SAS matmul library for Sparse Tensor Core. It is based on CuPy Okuta ct al. (2017) wrapper of cuSPARSELt. The index needs to be reordered before execution of sparse matmul (L8). Refer supp. E for more details about the reorder specific NVIDIA GPUs.
Listing 1: cuSAS:
The computational cost for the index is negligible (fig- report its wall-clock time); it merely looks at the sign bit of the log, neighbor activation (fig. 2p. Note: the SAS using this projection includes the ReLU network as a particular case (e.g., when odd elements of weight are zero for M-2). Other than this simple indexing, there are various options; one could choose N elements instead of 1 or incorporate another strategy for computing the index, c.g., estimating the index I using the previous layer's output (soction 6.1
FLOPS. The size of weight W for the SAS network increases linearly with M; nevertheless, the number of mult count is the same as the base dense/ /narrow network, i.c., C₁ x C, x HW. The mult count is also approximately the same with the SWS network (Zhou et al. 2021) with VM times wider input/output channel for 1:M sparsity pattern (Reler to supp. B for more detail). Both SAS and SWS require additional costs for routing the weight or activation. Sporse Tensor Core is highly engineered for this purpose and can execute the routing in parallel (section 2.3 and as reported in (NVIDIA) 2020), the time for the routing is small compared with the main matmul. The SAS network requires another negligible cost.
for computing index I of eq- 3 (fig. 3 reports the wall-clock time on actual hardware).
Memory. Although the mult count is constant for the activation sparsity M, the memory footprint of SAS increases linearly. This contrasts with the SWS, where storage for weight at inference time is constant irrespective of weight sparsity M. Note: the SWS network also requires M times more memory during training. SAS and the SWS require additional storage for storing the index; however, it is small, requiring only Mitt per element.
2.3 HARDWARE IMPLEMENTATION AND SPEED BENCHMARKING
The SAS matmul could be vectorized on an ordinary GPU having a local routing mechanism. For example, NVIDIA GPU equips this mechanism. During inference, we do not explicitly construct the sparse activation X; instead, we directly compute index I from dense acti- vation X (X is computed implicitly). The I and X is transferred to the core, then the
4

---

Published as a conference paper at ICLR 2024
Dense in CUBLAS (TF32)
500
CUSPARSELT
SAS in CUSPARSEL: (TF32)
400
Time (mes)
300
200
100
15360 14050 11520
2-3 size, y 10240
Figure 3: Speed benchmarking. SAS (1:2) vs SWS (1:2) for general matmul on NVIDIA A6000 GPU. Comparing the general matmul WX where X R and W R The extra cost for SAS, c.f., index computation is less than 1.5% (when M = 20480) of the entire matmul. The time of cuBLAS is an estimation based on NVIDIA's 's white paper and included for reference.
vectorized switch routes the corresponding weight element using the local index, executing the sparse/wide matmul as dense/narrow matmul(fig.) 2 in parallel. We developed general matmul library for SAS (listing [ called cuSAS for Sparse Tensor Core. cuSAS realized the vectorized matmul with structured sparse activation on GPU for the first time. Figure report the wall-clock time of SAS for general matmul compared with SWS using the same matrix size. The difference between SWS and SAS comes from the overhead specific to SAS (index computation, reorder (supp. E and its memory transfer), which is less than 1.5%, even with naive implementation having several redundancies. Note: the configuration of this experiment is different. from the previous discussion and the other experiments using neural networks; we use the same matrix size for dense, SWS, and SAS to reveal their dif- ference more clearly. Therefore, mult count of SWS and SAS is 1/M of that of dense when considering the sparsity. Refer to supp. C for more detail about the speed benchmarking.
2.4 TRAINING SAS NETWORK
One could transform most existing (narrow) neural networks into (wide) SAS networks to increase their expressive power without increasing inference speed. The conversion is simple: replacing the original non-lincarity (c.g., ReLU) with our SAS projection S. During the training, SAS matmul is computed as follows: it first projects the dense/narron nnarrow activation map inlo structurally sparse/ wide space by S, constructing the sparse activation explicitly; then, it performns the conventional dense matmul. It is equivalent to the sparse matmul for efficient inference using hardware support (fig.
Optimizer for SAS To train the SAS network, one wants to utilize modern optimizers such as Adam (Kingna & Ba 2014) and AdamnW (Loshchilov & Hutter 2019). which use adaptive learning rates based on the statistics of the gradient history. However, it is known that the adaptive learning rate could be unstable in the early stage of training, where the number of experienced gradients is small (Liu et al. 2020). This is more problematic in the SAS network (when M is large) because each weight element receives gradients less frequently (x (x1/M on average) than the dense network. To achieve stable training of the SAS AnperienceRRAdore (ERAdam), which slightly modifies the formulation of RAdam (Liu et al. 2020). Basically, RAdam rectifies Adam W)'s adaptive learning rates by multiplying r < to reduce the variance in the early stage of training:
network, we propose
Poo 2 21 (1) 32) 1, Pt = Px 21/8 (1-13)-
(4)
Ti ((x (p: 2) pr)) 10))))) (/(x 4) (prx - 2) pt).
(5)
where t is the optimization steps, and B is the hyperparamcters used in computing the running average of gradients. As + oo (training procoods), we have r. - 1, and it recovers the original Adam(W)'s update.
Our ERAdam scales steps for calculating 7, of each weight element so that it is proportional to the number of received gradients. Specifically, we the following scaled time vector
5

---

t e RCC instead of Published as a conference paper at ICLR 2024
Activarion Sparse
Input (1=1.0)
Rell (t=0.89)
M=2 (1=1.2)
M=4 (183.8
M= (t=5.2)
M=16 (1=6.3)
Weight Sparse
Input (t=1.0)
BeLU (l=0.89)
M=2 (I=0.99) M=4 (I=0.99)
M=8 (I=0.99) M=16 (1=0.99)
Figure 4: SAS capacity analysis by trajectory length. Comparing SAS (top) with dense ReLU
network and SWS network (bottom). indicates the average relative length concerning the input circle length, where longer length (complex trajectory) indicates expressive power.
t++ ((VXX..) #0) /HW,
(6)
where VX E C, HW is gradient activation corresponding to the weight W. We find the optimizer with the CosineDecay (Loshchilov & Huiler 2017) scheduler with k-decay ZZhan & 2020 works well, especially for large M.
EXPRESSIVE POWER ANALYSIS
We propose SAS for DNNs, but it is difficult discuss its superiority (or inferiority) over SWS based on their accuracy on a specific task; the result often turns back depending on the model structure, training strategy (scheduler, optimizer, etc), and dataset. To evaluate the difference of expressive power with SWS independently from these factors, we utilize the Trajectory Length (Raghu et al.) (017)). which evaluates the expressiveness by measuring the length of the output trajectory given input sweeps along a one-dim path (e.g., circle).
We construct a two-layer neural network with 2-dim inpul and output. (fig, A1). The base network has 32-dim di- intermediate channels, followed by the ReLU activation. We replace the ReLU by our SAS with different sparsity M, such that they all have the same mult count, i.e., structurally sparsified activation X has nonzero elements. We also construct the SWS variant (Zhou et al. 2021). which has approximately the same mult count (supp. section A). We evaluate their trajectory length by inputting the one-dimensional circular path the networks by varying the sparsity M and report their average oulput trajectory length by changing the initial random weights.
Figure repoort the results. The trajectory length of the SAS net increases with increas- ing sparsity M. On the other hand, the length from the SWS network almost constant for M. This trajectory analysis suggests SAS has more expressive power than the SWS, given the same computational budget at the expense of increased memory requirement.
EXPERIMENT
Given the same ma natmul budget, enlarging model by SAS with sporse projec- tion improves accuracy? Is il better than the SWS using the same hardware?
We conducted extensive experiments to answer the questions. For the first question, we evaluated the accuracy by changing the sparsity M, keeping the same mult count. For the second question, we compared the SAS with SWS (Zhou et al., 2021). which has the same sparsity and mult count; both utilize the same hardware mechanism (i.e., local routing).

---

Published as a conference paper at ICLR 2024
CIFAR--1
0.4
Train summaar
Train, 16
Train,
Train,
Test summary
Test, 16
Test,
Test, C
0.5
CIPAR-100
0.
Train summary
Train, 16
Train,
Train,
3/5
3/3
Test, Figure 5: Accuracy of narrow ResNet18 CIFAR-10 and CIFAR-100 for difference SAS
Test summary
Test, 16
Test, a=
sparsity M. The summary plot reports the best score for each setting where dashed lines are the SWS network having the same sparsity weight, which consumes the same mult-add count.
EXPERIMENTAL SETUP AND RESULTS
We evaluate the classification accuracy CIFAR-10/100 (Krizhevsky et al. b), and ImageNet (Deng et al., 2009). Because the commonly used network already achieves good saturated accuracy, we adopted the narrow version as a base network better evaluate the effects of using the sparse activation or weight. Specifically, as base dense ReLU network, we've utilized the network with ax narrow width (W CRCo/xCC/o)) except for the first and last layers; it approximately consume a²x less mult count. All variants, ReLU, SWS (Zhou al.) 2021 and SAS networks, have approximately the same mult count for the same The I:M SWS network employs vM* x more channel to make the mult count the
same as the base ReLU network (section 2.2 supp. B
For fair comparisons, employ the same training strategy for all variants, ie., the training epochs, batch size, optimizer, scheduler, etc. For training the SWS network, we employed the training method of (Zhou et al.) 2021). We train the networks from scratch for all experiments. Refer to supp. D for more dotail about the experimental configuration.
CIFAR-10 CIFAR-100. We use ResNet18 (He al. 2016) architectures. For all the variants, we utilize our proposed ERAdam optimizer (section 2.4] in combination with k-decay scheduler (Zhang & Li 2020); it is equivalent to RAdam for
of the most popular

---

Published as a conference paper at ICLR 2024
the base network because En( (VX.,, + in eq. (6 is approximately same for each weight elements. We trained them for 16x 1000/a epochs, utilizing the data argumentation adopted in the ConvNeXt (Liu et al. 2022). We evaluate the accuracy by changing the base network summarize the results.
width a and sparsity M. Figure
ImageNet. We use ConvNeXt (Liu et al. 2022), which is computationally efficient by its network architecture design; c-g., it heavily utilizes dopthwise convolutions to reduce the number of mult. For this ImageNet experiment, we conducted an experiment only on the practical configuration (rather than changing and M in wider rage); we use ConvNeXt-B, having half of the original input channel (a = 2) as a base network and use M 2 for SWS and SAS. We use their official code base only changing its activation and number of training epochs and keeping the rest of the highly optimized settings unchanged. The accuracy of the base network, SWS, and SAS is 80.7 %, 81.5 %, and 82.2 %, respectively.
4.2 DISCUSSION
As expected, the accuracy improves as we increase the activation's sparsity M. For all ranges of the base net width (having different a), we observed the most significant gain from the base ReLU network to M-2 sparse network. The improvement saturates around M=8 M-16. Similarly,
observe significant gain when the base network capacity is small (large a). From a perspective of the training statistics evolution, the accuracy of the SAS network having significant sparsity (large M) is worse than the less sparse network (small M) at the beginning of the training; however, it surpasses them as training procoeds; sparser network requires more training epochs to reach their best accuracy. We hypothesize this is partly due to the sparse gradient when M is large. In comparison with the weight sparsification, SWS network (Zhou et al., 2021) also shows improvement for the weight sparsity M. Yet, SAS scales significantly better than SWS for the same mult count.
In summary, we see the SAS realized by the projected sparsity boosts accuracy without increasing mult count; furthermore, the improvement is monotonic w.r.t. M and is much better than SWS. The index I changes depending on its input, suggesting the different weight elements are dynamically utilized depending on their imput (section 5). We hypothesize that the increased flexibility by the increases memory for weight a primary source of SAS's superior accuracy boost over the SWS. The result also aligus with the increased capacily of SAS evaluated by the trajectory length analysis (section
RELATED WORKS
This section discusses the relation other approaches for efficient DNN inference GPU.
Unstructured weight sparsity. It is hard to utilize the unstructured sparsity for speedup on GPU (Shi et al.] 2020). For example, the wall-clock time for matmul be- tween an 8000x8000 matrix with sparsity of 90% and a dense matrix with the same size lakes 780ms by cuSPARSE, while the corresponding dense algorithm by cuBLAS only requires 121ms. Though the sparse operation reduces the mult count by 90%, the dense matmul is 7x faster, suggesting the difficulty of utilizing unstructured sparsity in vector processors due to several overheads such as indexing.
Dynamic kernel. SAS dynamically selocts weight elements basod on activation. In this regard, it relates to series of rescarch exploring the input-dependent weight to increase the network expressive power (Chen al. 2020 Jia et al. 2016)., which also relates to the attention mechanism /Vaswani et al. 2017 Our SAS is different both in its molivat ion and
mechanism (e.g., SAS selects weight instead of composing it on the fly). Possible future work would be to use attention for computing the index instead of the strategy in section 2.2
Quantization. Discretizing the weight and/or activation into a low-bit representation is commonly used to speed up the DNN inference (Yin et al. 2019 Esser et al.) 2019 Li & Baillieul 2004); low-precision mult is more computationally efficient and consumes less energy than the floating point counterpart. Furthremore, when we adopt lower bit weight,
8

---

Published as a conference paper at ICLR 2024
memory footprints and cost for transferring the weight decrease. The main drawback of SAS the increased memory footprint for the weight matrix (xM); we'll explore the combination with the lower-bit weight matrix to mitigate the issue.
Low-rank factorization. Factorization of matmul into a low-rank presentation is often utilized to reduce the number of mult count. The factorization can be integrated into network design as depthwise or Ixl convolution; they are utilized in efficient network designs such as MobileNet (Howard et al., 2017) or ConwNeXt (which we evaluate). Recent research realized the optimization of low-rank structure (Idelbayev & Carreira-Perpinat Perpinan 2020 Ours are orthogonal these low-rank approaches and could be combined.
CONCLUSION
We propose SAS, an unexplored but effective approach for efficient DNN inference on vector- type processors. The evaluation reveals that sparsifying activation by structured projection improves accuracy while maintaining the same mult count and inference speed. Our idea is compatible with NVIDIA's commercial GPU for M-2. It is not limited to the specific device but good match for a wide range of vector processes; we expect door for the new algorithm and hardware utilizing the structured sparsity in activation.
open the
6.1 LIMITATIONS & FUTURE WORK
Advanced projection algorithm. In this study, we employed a straightforward strategy to compute the index I (section 2.2 We employed this method because it is computation- ally cheap, simply looking at the sign bit of the neighbor activation. It has some trawbacks for M>2 when the neighbor elements are close to zero. When the neighbor oscillates around zero, the sparse activation undergoes discontinuous changes, which may negatively affect the accuracy; this phenomenon is more likely to occur when M is large. From another perspective, our current implementation determines the index by the simple rule, which is not learned end-to-end. Learning to compute the index by incorporating extra. channels or attention-like mechanisms (Vaswani et al. 2017) might improve the accuracy.
Combination with SWS. We used the switching circuit developed for the SWS utilize the sparsity in activation for speedup. Exploring the combination of SAS and SWS would be an exciting direction. Unfortunately, the current GPU supports sparse matmul where one of the matrices is sparse. More efficient DNN could be realized on future GPU, which supports matmul where both inputs are structurally sparse.
Library for DNN framework. We developed the CuPy Okula et al. 2017 based li- brary cuSAS (listing I for general matmul for 1:2 sparse pattern, which runs on actual Sparse Tensor Core hardware and realized wall clock speedup; however, integration to pop- ular DNN framework has not yet been completed. As an important future work, we will develop SAS (depth-wise) convolution for Pytorch (Paszke al.) 2019), TensorFlow (Abadi ct al.) 2015). and JAX (Bradbury et al.) 2018) by CUDA programming. We expect slightly better results (in terms of the percentage of SAS-specific overhead) than the result in fig. will be achieved with the custom CUDA kernel because we can integrate SAS-specific online computation such as index computation, reorder (section E and memory copy.
Hardware consideration for significant. sparsity. Current NVIDIA's GPU support only 1:2 sparsity for tf32 or float matricos which restricts the SAS in higher sparsity on actual hardware. In principle, larger sparsity in SAS does not increase the mult count but may increase the switching circuit's area, energy, and storage cost. Considering the saturated gain by increasing the sparsity (section 4 the best sparsity can be defined; we left this consideration for future work. Furthermore, hardware support could make the index computation (section 2.2) more efficient. In addition to the efficient implementation on existing GPUs, we'll explore the novel vector-type processor design specific to the SAS.
https://docs... nvidia. wwwdia..cm//uuaaaaaaaa. com/cuuda/cuepprreltt Tt supports 2:4 sparsity for lower hit inputs.

---

Published as a conference paper at ICLR 2024
REFERENCES
Martin Abadi, Ashish Agarwal, Paul Barham, Eugene Brevdo, Zhifeng Chen, Craig Citro, Greg S. Corrado, Andy Davis, Jeffrey Dean, Matthieu Devin, Sanjay Ghemawat, lan Goodfellow, Andrew Harp, Geoffrey Irving, Michael Isard, Yangqing Jia, Rafal Juzefowicz, Lukasz Kaiser, Manjunath Kudlur, Josh Levenberg, Dan Mané, Rajat Monga, Sherry Moore, Derek Murray, Chris Olah, Mike Schuster, Jonathon Shlens, Benoit Steiner, Ilya Sutskever, Kunal Talwar, Paul Tucker, Vincent Vanhoucke, Vijay Vasudevan, Fernanda Viégas, Oriol Vinyals, Pete Warden, Martin Wattenberg, Martin Wicke, Yuan Yu, and Xiaoqiang Zheng. TensorFFoow Large-scale machine learning on heterogeneous systems, 2015. URL http://tensorflow org/ Software available from tensorflow.org
Jorge Albericio, Patrick Judd, Tayler Hetherington, Tor Aamodt, Natalie Euright Jerger, and Andreas Moshovos. Cuvlutin: Irefo deep neural network computing. In 2016 ACM/IEEE 43rd Annual International Symposium on Computer Architecture (ISCA), pp. 1-13, 2016. doi: D0.1109/ISCA.2016..0
James Bradbury, Roy Frostig, Peter Hawkins, Matthew James Johnson, Chris Leary, Dou- gal Maclaurin, George Necula, Adam Paszke, Jake VanderPlas, Skye Wandernan-Milne, and Qiao Zhang. JAX: composable transformations of Python+Numll programs, 2018. URL
http://ittubb.ccoooogo...e/
Tianlong Chen, Xuxi Chen, Xiaolong Ma, Yanzhi Wang, and Zhangyang Wang. Coarsening the gramularity: Towards structurally sparse lottery tickets. February 2022.
Yinpeng Chen, Xiyang Dai, Mengehen Liu, Dongdong Chen, Lu Yuan, and Zicheng Liu. Dynamic convolution: Attention over convolution kernels. In Proceedings of the IEEE/CVF conference on computer vision and pattern recognition, pp. 11030-11039, 2020.
Deng, Wei Dong, Richard Socher, Li-Jia Li, Kai LI, and Fei-Fei. ImageNet: large- scale hierarchical image database. In 2009 IEEE Conference on Computer Vision and Pattern Recognition. IEEE, June 2009.
Steven K Esser, Jeffrey L McKinstry, Deepika Bablani, Rathinakumar Appuswany, and Dharmen- dra S Modha. Learned step size quantization. February 2019.
Jonathan Frankle and Michael Carbin. The lottery ticket hypothesis: Finding sparse, trainable neural networks. March 2018.
G. Georgiadis. Accelerating convolutional neural networks via activation map compression. In 2019 IEEE/CVF Conference on Computer Vision and Pattern Recognátion (CVPR), pp. 7078 7088, Los Alamitos, CA, USA, jun 2019. IEEE Computer Society. doi: N0..109/CVPR.2019..00 URL hhttp::///oo wtw.::eeecccccppcscseeeeeeeeeeee www.iiettt..or.eeeeeeoe 0t///00.11000 11w.//VPR./C...... 2019 2019.00725
Song Han, Xingyu Liu, Huizi Mao, Jing Pu, Ardavan Pedram, Mark Horowitz, and William Dally. EIE: Efficient inference engine on compressed deep neural network. February 2016.
Kaiming He, Xiangyu Zhang, Shaoqing Ren, and Jian Sun. Delving deep into rectifiers: Surpassing human-level performance on imagenet classification. In Proceedings of the IEEE international conference on computer vision, pp. 1026-1034, 2015.
Kaining He, Xiangyu Zhang, Shaoqing Ren, and Jian Sun. Deep residual learning for image recognition. In IEEE Conference on Computer Vision and Pattern Recognation, pp. 770 778, 2016.
Geoffrey Hinton, Oriol Vinyals, and Jeff Dean. Distilling the knowledge in a neural network. orXiv preprint or arXiv:1533.0255 2015.
Andrew Howard, Menglong Zhu, Bo Chen, Dmitry Kalenichenko, Weijun Wang, Tobias Weyand, Marco Andreetto, and Hartwig Adam. MobileNets: Efficient convolutional neural networks for mobile vision applications. April 2017.
Yerlan Idelbayev and Miguel A Carreira-Perpinan. Low-rank compression of neural nets: Learn- ing the rank of each layer. In 2020 IEEE/CVF Conference on Compuler Vision and Pallern Recogriition (CVPR). IEEE, June 2020.
Xu Jia, Bert De Brabandere, Tinne Tuytelaars, and Luc V Gool. Dynamic filter networks. Advances in neural information processing systems, 29, 2016.
10

---

Published as a conference paper at ICLR 2024
Diederik P Kingma and Jimmy Ba. Adam: A method for stochastic oplimization. December 2014.
Alex Krizhevsky, Vinod Nair, and Geoffrey Hinton. Cifar-10 (canadian institute for advanced research). URT. httt:://ww/ cs. www.oo.tttoooooooooo serronto.edu/krriie edu/-kriz//ifarr html
Alex Krizhevsky, Vinod Nair, and Geoffrey Hinton. Cifar- 100 (canadian institute for advanced research). b. URL htttp://ww..cs torontoo wot.oooo...ooorrraaaat.. www//krii/ http://iarr.../ cifar.html
Mark Kurtz, Justin Kopinsky, Rati Gelashvill, Alexander Maiveev, John Carr, Michael Goin, William Leiserson, Sage Moore, Nir Shavit, and Dan Alistarh. Inducing and exploiting activation sparsity for fast inference on deep neural networks. 119:5533 5543, 2020.
Hao Li, Asim Kadav, Igor Durdanovic, Hanan Samet, and Hans Peter Graf. Pruning filters for efficient convnets. arXiv preprint arXiv:608..000 2016.
Keyong Li and John Baillieul. Robust quantization for digital finite communication bandwidth (dfeb) control. IEEE Transactions on Automatic Control, 49(9):1573-155) 1584, 2004.
Y. Li, K. Adamezewski, W. Li, S. Gu, R. Timofte, and L. Van Gool. Revisiting random channel pruning for neural network compression. In 2022 TEEE/CVF Conference on Computer Vision and Pattern Recognation (CVPR), pp. 191 201, Los Alamitos, CA, USA, jum 2022. IEEE Computer 202.PP5268.......002 2022.00029
Society. doi: 10.1109 109/CVPR5268S
URL https://doi ieeecomputersocietyycomm org/
10. 110.9//VPPPS288.... 2022. .00029
Liyuan Liu, Haoming Jiang, Pengcheng He, Weizhu Chen, Xiaodong Liu, Jianfeng Gao, and Jiawei Han. On the variance of the adaptive learning rate and beyond. In International Conference on Learning Representations, 2020. URL
Zhnang Liu, Hanzi Mao, Chau- Wu, Christoph Feichtenhofer, Trevor Darrell, and Saining Xie. A convnet for the 2020s. In Proceedings of the IEEE/CVP Conference on Computer Vision and Pattern Recognition, pp. 11976611986, 2022.
Tlya Loshchilov and Frank Hutter. SGDR: Stochastic gradient descent with warm restarts. In International Conference Learning Representations, 2017. URL https:///ppareviev..ett forum?id forum?id=SSqg9Scez
Tlya Loshchilov and Frank Hutter. Decoupled weight decay regularization. In International Confer- ence on Learning Representations, 2019. htt.:g/gg//pppaaaaaaaaaa/a//e/er
NVIDIA. NVIDIA A100 tensor core GPU architecture. httpps://inages.nvidae hww.://..iiidiiaaccmmmmmeaaaaa het///Slltoons/dattt. center/nvidia amplra-aampreeeee amppre-archatectuueeeee archhhattctccteettttttpppprrpppa 2020.
Ryosuke Okuta, Yuya Unno, Daisuke Nishino, Shohei Hido, and Crissman Loomis. Cupy: A numpy- compatible library for nvidia gpu calculations. In Proceedings of Workshop on Machine Learning Systems (LearningSys) in The Thirly-first Annual Conference Neural In Information Processing Systems (NIPS), 2017. URL http:
Angshuman Parashar, Minsoo Rhu, Anurag Mukkara, Antonio Puglielli, Rangharajan Venkatesan, Brucek Khailany, Joel Emer, Stephen W. Keekler, and William J. Dally. Scnn: An accelerator for compressed-spare sparse convolutional neural networks. In Proceedings of the 44th Annual Interna- lional Symposium on Compuler Archsilecture, ISCA '17, pp. 27a 40, New York, NY, USA, 2017. Association for Compuling Machinery. ISBN 9781450348998.. doi: 00.11453379856.333 URL https://doi. org/ /00 01445/079856... 307985663080254
Jongsoo Park, Sheng Lı, Wei Wen, Ping Tak Peter Tang, Hai Li, Yiran Chen, and Pradeep Dubey. Faster CNNs with direct sparse convolutions and guided pruning. In International Conference Learning Represenlations, 2017. URL httpss//openreviei wwwparreiieeeee/oouuueeeeeeeeeee
Adam Paszke, Sam Gross, Francisco Massa, Adam Lerer, James Bradbury, Gregory Chanan, Trevor Killeen, Zeming Lin, Natalia Gimelshein, Luca Antiga, Alban Desmaison, Andreas Kopf, Edward Yang, Zachary DeVito, Martin Raison, Alykhan Tejani, Sasank Chilamkurthy, Benoit Steiner, Tu Fang, Junjie Bai, and Soumith Chintala. Pytorch: An imperative style, high-performance deep learning library. In H. Wallach, H. Larochelle, A. Beygekzimer, F. d' d'Alcé-BBue, E. Fox, and R. Garnelt (eds.), Advances in Neural Information Processing Systems 32, pp. Curran Associates, Inc., 2019.
Maithra Raghu, Ben Poole, Jon Kleinberg, Surya Ganguli, and Jascha Sohl-Dickstein On the expressive power of deep neural networks. In international conference machine learning, pp- 2847 2874. PMLR, 2017.
11

---

Published as a conference paper at ICLR 2024
Minsoo Rhu, Mike O'Connor, Niladrish Chatterjee, Jeff Pool, Youngeun Kwon, and Stephen W. Keckler. Compressing dma engine: Leveraging activation sparsity for training deep neural net- works. 2018 IEEE International Symposium on High Performance Computer Architecture
(HPCA), pp. 78-91, 2018. doi: 10.1109/HPCA.2015 SAAPCCA.011.00000 00017.
Shaohuai Shi, Qiang Wang, and Xiaowen Chu. Efficient sparse-dense matrix-matrix multiplication on gpus using the customized sparse storage format. In 2020 JEEE 26th International Conference Parallel and Distributed Systems (ICPADS), pp. 19 26. IEEE, 2020.
Zhanhong Tan, Jiebo Song, Xiaolong Ma, Huat Tan, Hongyang Chen, Yuanqing Miao, Yifu Wu, Shaokai Ye, Yanzhi Wang, Dehui Li, et al. Penn: Pattern-based fine-grained regular pruning towards optimizing enn accelerators. In 2020 57th ACM/TEEE Design Automation Conference (DAC), pp. 1-6. TEEE, 2020.
Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N Gomez, Lukasz Kaiser, and Illia Polosukhin. Attention is you need. Advances in neural information processing systems, 30, 2017.
Yang Wang, Chen Zhang, Zhiqiang Xie, Cong Guo, Yunxin Liu, and Jingwen Leng. Dual-side sparse lensor core. In 2021 ACM/IEEE 48th Annual International Symposium on Compuler Architecture (ISCA), pp. 1083 1095. IEEE, 2021.
Yanzhi Wang, Shaokai Ye, Zhezhi He, Xiaolong Ma, Linfeng Zhang, Sheng Lin, Geng Yuan, Sia Huat Tan, Zhengang Li, Deliang Fan, et al. Non-structured dmn weight pruning consid- ered harmful. Xii preprint arXie:1907.02123 2, 2019.
Wei Wen, Chunpeng Wu, Yandan Wang, Yiran Chen, and Hai Li. Learning structured sparsity in deep neural networks. Advances in neural information processing systems, 29, 2016.
Mitchell Wortsman, Ali Farhadi, and Mohammad Rastegari. Discovering neural wirings. June 2019.
Huanrui Yang, Wei Wen, and Hai Li. DeepHoyer: Learning sparser neural network with differen-
tiable Scale-Invariant sparsity measures. September 2019.
Penghang Yin, Jiancheng Lyu, Shuai Zhang, Stanley Osher, Yingyong Qi, and Jack Xin. Under- standing Straight-Throug estimator in training activation quantized neural nets. March 2019.
Tao Zhang and Wei Li. kdecay: Just adding k-decay items on learning-rate schedule to improve neural networks. arXio preprint arXi::2004.0500 055909 2020.
Aojun Zhou, Yukun Ma, Junnan Zhu, Jianbo Liu, Zhijie Zhang, Kun Yuan, Wenxiu Sun, and Hong- sheng Li. Learning N:M fine-grained structured sparse neural networks from scratch. February 2021.
Zeqi Zhu, Arash Pourtaherian, Lue Waeijen, Lennart Bamberg, Egor Bondarev, and Orlando Mor- eira. Arts: An adaptive regularization training schedule for activation sparsity exploration. In 2022 25th Euromieero Conference on Digital System Design (DSD), pp. 415 422, 2022. doi: 2.1100/DDD577022.2222.0000
12

---

Published as a conference paper at ICLR 2024
SAS: Structured Activation Sparsification Supplemental Material
A DETAIL IN TRAJECTORY LENGTH ANALYSIS
Network design for the Trajectory Length analysis (section 3 is illustrated infig. A1 We
evaluate the length by randomly initializing the weight 100 times and reporting their average. We also compose the SWS network having the sparsily, which consumes approximately the same mult count. The SWS network uses ReLU activation.
(5,9)
(8,9)
& 8 8 8
BAD
(weight sparse) (1:2)
(activation sperse)
Figure A1: Network design used for the Trajectory Length Analysis (section Relative output length with respect to the imput circle length is an indicator of the network expressive power; longer length (complex trajectory) indicates more expressive power.
B DETAIL IN SWS NETWORK
In this paper, we propose SAS to improve the network capacity or accuracy without increas- ing the actual mult count. Therefore, we compare the SWS in the same scenario (keeping the mult count constant while increasing the sparsity). Specifically, we consider a base layer consisting of matmul between activation X and weight W, where X e &&x and W E RCxxC; (eq. (1). Using the proposed SAS, one could increase the network width for M times while maintaining the same mult count as the base layer by utilizing the 1:M sparsity pattern. The sparsified matrix shape is, R MC, and W &C.xMCC (eq. (3) Note that the SAS does not change the output channel dimension Co.
On the other hand, in the case of SWS, if one increases the network width for the M times and uses the 1:M sparse pattern on weight, the mult count of the resultant network increases by about M x M/M = M because both input and output channel is M times wider. In the case SWS, by using VM times wider network for the 1:M sparsity pattern, we can construct the SWS network, which has roughly the same mult count the base dense network and SAS. More specifically, we adopt the following configuration: We use VM times (instead of M times) wider approximately the same mult count as the base network. The network width needs to be an integer value, and it also needs to be a multiple M. Hence, we use the weight having the shape of LVMC
channel. this configuration, we'll have
LVMC!!1 for the 1:M SWS network, where L-1 is rounding operator, (1)) (() is the input/output channel dimension of 1-th layer of the base dense (LVMC("[LMMC)
network. For the last chunk, when it does not equal to M,
13

---

Published as a conference paper at ICLR 2024
mod M):1 sparse pattern. This way, construct the SWS network having approximately the same mult count as the base dense network.
For example, consider the I-th layer of the base network consisting of convolution with C,=288 (C,=32, kernel size k=3) and C,=128. When M-8, then we have 1=80 and 1V1 Then mult count of the original dense layer (for single pixel) is 288x64=1843,, the mult count. of the weight sparse layer (SWS) is 18440 (815x
with modulo 3, we use M-3 for the last chunk).
C DETAIL SPEED BENCHMARKING
we adopt the opposite config- uration as the neural network experiment. We evaluate the speed by changing the sparsity pattern while keeping the matrix dimension unchanged to report more concrete comparisons.
On the wall-clock speed benchmarking reported in section 2.3
Specifically, the wall-clock speed benchmarking in fig, 3 we consider the matmul WX where X IR*x and W E Rx97 which is same for dense matmul, SAS matmul, and SWS matmul. The mult count the three variants the same x when one does not take the sparsity into account. By utilizing the 1:2 fine-gre ane-grained (or semi-structured) sparsity on Sparse Tensor Core, the mult count becomes half for SAS and SWS, i.e., (Y x a x B2/ fixed 7=10240, changing = 3 from 10240 to 20480, which is
As shown in fig. the same for dense matmul SAS matmul, and SWS matmul.
We want emphasize that the scenario in this speed benchmarking (keep the same matrix dimension) is different from the scenario for neural networks (keep the (almost) same mult count). We adopt the different configurations to evaluate the speed in fair setting between SWS and SAS. In the neural network setting, one could construct the SWS network having approximately the same mult count with the base one by using VM wider input/ /out/uu channel (supp. B); however, it is hard align the mult count preciscly, and the shape of the matrices is different, which makes it hard to evaluate the overhead specific to SAS which we are interested in (computation of index, reorder, and memory transfer for the index). On the other hand, we can clearly evaluale this by measuring the time of the matmul of the same-sized matrix (section 2.3
14

---

Published as a conference paper at ICLR 2024
D DETAIL IN MAIN EXPERIMENT
The primary experimental setup is summarized in table A1 The FLOPS and memory
footprint of the network used in the experiments are summarized in table A2
For training the SWS network, we employed the method of (Zhou et al 2021 instead of
APEX's Automatic SParsity beccsss ily 1: : sparsity and @2it allows training from scratch which enables fair comparison with SAS.
because Ocode base of (Zhou et al. 2021
supports arbitrar-
Table A1: Experimental setup
CIFAR-10
CIFAR-100
ImageNet
Network
ResNet18
ConvNeXi-B
Batch size
512
4096
Training epachs
16/ax1000
600
Optimizer
ERAdam (section
Adamm
Scheduler
Two cycle cosine with kDecay=2.0 Zhang & Li 2020
Cosine
Initialization
Kaiming-umiform uniform He al. 2015)
Truncaled Gaussian
Base width
4/8/16
2
Sparsity
ReLU/2/4/8/16
ReLU/2
CIPAR-10/CIPARC
The code for CIFAR-10 and CIFAR-100 is based on Pytorch lightning CIFAR10 tutorial code 3 We use a single A6000 GPU for CIFAR10 and CIFAR100 experiments. Tt. takes a day to train the single model.
IMAGENET
The code for ImageNet is based on Conv (Liu et al. 2022) four A100 GPUs (each holding 256 batches) with an update frequency of four to virtually construct the batch size of 4096. It takes about two weeks to train a single SAS model for 600 epochs (double the original 300 epochs to compensate for sparse gradient); we use the default setup of ConvNeXt's s (Liu et al. 2022) official repository for training ImageNet-1k (without pre-training using ImageNet-2224 22K), only changing the original dense matmul to our proposed SAS matmul (con - SASconv2D, and linear SASlinear) and training epochs (300 - 600). We use the original AdamW optimizer to keep the original ConvNeXt's highly optimized settings intact much possible (furthermore, in this moderate sparsity of M=2, proposed ERAdam behaves almost identical to AdamW with warm-up). Refer to their paper for a more detailed setup.
official code base We use
4
Table A2: FLOPS and memory footprints. Note that FLOPS and the number of parameters and FLOPS of SWS are not preciscly the same the base dense network as discussed in supp.
B but the difference is less than 1% for all the configurations.
Network
FLOPS (all)
Params (dense, SWS)
Params (SAS)
ResNet18 (a=4)
114K
731K
x31KxM
ResNet18 (a=8)
28K
182K
182K X
ResNet18 (a=16)
7K
46K
46K x M
ConvNeXt- (a=2)
3850M
22M
22Mx
http://atthbb.co///DDD///ap
http: http://iaaaag..ooooprrrcooooooor
15

---

Published as a conference paper at ICLR 2024
TRAINING
During the training, SAS matmul is computed as follows: it first projects the dense/narrow activation map into a structurally sparse/we heaars/wwdeaaae space by S, constructing the sparse activation explicitly; then, it performs the sonventional dense matmul. Note that it is equivalent to the sparse matmul for efficient inference when hardware support is available, and do not construct the sparse matrix explicitly during inference. Still, the hardware directly processes the narrow dense feature aloug with the index computed online (fig, 2 Refer to the pscudo-code in listing 2 during training. When implemented this way, the gradient for the wide weight W could be computed using autograd mechanism.
del SAS.
corresponds log2(M) the main x.shape
3 xa (torch.roll(xf: (i Nonel, i, dims=1) for in range(m)] }) for in onumerate[sa]], dim=2)).ummdiun (amimmmmiii keepdin - True) x_sparse terch 2'm. H. Wi).scatter__Q (1, ind, xt. viiew). (2'm)+C, H, W])
4 ind torch.cat Catch.cat(2eeeeeeee
return -sparae
def forward(sell, x): # x: imput activation, mi sparsity factor (actual sparsity is 100(1-111)) 1000-1/2200 1/22")]/))
return F.com22d(SA8 proj(x sellwwight, biss=salf.bias, stride=....)
Listing 2:
Code of SAS conv2d layer for training (PyTorch). Note: During inference, does not need construct the sparse activation explicitly (L5); refer fig. for efficient infrence mechanism.
16

---

Published as a conference paper at ICLR 2024
E MEMORY ARRANGEMENT FOR CUSPARSELTMATMUL
The Sparse Tensor Core and the cuSPARSELt library were originally developed to speed up the DNN having structured weight sparsity (SWS). Our SAS improved the accuracy cetturacy/cooppii /compu- tation tradeoff by using the same hardware and software library. In the case of SWS, the index could be precomputed after training using cusparsel.&pp/WACCpppee. function of cuSPARSELt library. Figure A2 illustrates the memory arrangement of value and index for cusparseLtMatmul operation for NVIDIA's Sparse Tensor Core. The value matrix follows the index matrix. In the case of a 1:2 sparse pattern (TF32), the index is stored using 4-bit (although it can be represented in 1-bit; Sparse Tensor Coreuse 4-bit to index a 1:2 pattern), and '0x4' and 'Oxe' is assigned for indexing the first and second element, respectively.
The important note for Sparse Tensor Core is that the order of the index not aligned with its corresponding activation value; it needs to be reordered (L8 in listing 1 and the reordered index I needs to be supplied to the core to get the correct result. The arrangement depends on the size activation matrices X (The memory arrangement in fig. A2 illustrates the specific case when the input matrix is 64 32).
HOW TO USE THE CUSPARSELTMATHUL FOR SAS?
One can also use the cusparselttppMAConpeerss function for SWS to executed the SAS matmul, Ccompute the index I using equation 3 then @explicitly computing the sparse activation X and finally @compute the reordered index I using cusparseLttppMt.Cppeess. However, it is redundant and inefficient. We already have compressed activation X as an output from the previous layer; the index I is computed cheaply. We want to reorder the index I to get I without explicit construction of sparse activation X discussed in the main text (section 2.3 The problem is that NVTDIA does not provide information about the rendering mechanism of CuSPARSELt.
TOOLS FOR DIGGING
REORDERING MATRIX
We developed a helper tool of cuSAS for finding out the rendering matrix O for arbitrary size matrix realize general matmul for SAS activation. The core idea for realizing the elucidation is impulse response of cusparssalt.ppMacoprres function. Specifically, we input the sparse activation X; all the odd element has a non-zero value except the (i.j) element. Then, looking at the reordered index compuled by cusparseeltSpMMMCopprrrr we can find the destination index (i.5) where (i.j) element the original index should be warped. Repeating this process for all the row-column pairs, get the reordering matrix O such that I = Z[O). The reordering matrix O depends only on the size of X; therefore could be precomputed. We'll also open-source this tool.
It is possible to implement the CUDA kernel, which runs the following operation at once for more efficient SAS matmul; Dcheeks the sign bit, Dassign either '0x4' or 'Oxe', and Dreorder.
M
M x K x:32
Value
Index
Figure A2: Memory arrangement for cusparseLtNatmil. The index is located just after the
value. The index needs to be arranged for the execution of sparse matmul using cusparselcMatmul. The formatting of the index depends on the size of the matrix in listing [1] This figure illustrate the case for the 1:2 sparse pattern (TF32) and the input matrix size 32 64.
TThe current version supports sparse activation with 1:2 structured pattern.
17

---

Published as a conference paper at ICLR 2024
F FINE-GRAINED (SEMI) STRUCTURED WEIGHT SPARSITY.
NVIDIA 2020 could speed up the matmul with weight having moderate rate sparsity (e.g.. 50%) GPU, which is almost impossible for the unstructured sparse pattern (section 5 They realized actual speed up by utilizing a specific pattern in their sparsity, namely N M structured sparsity. Suppose typical matrix multiplication between activation X E R16x32 and weight W E R232 The Dense Tensor Cores implement this matmul by two cycles. In contrast, the Sparse Tensor Core only needs one cycle if the weight tensor W satisfies the structured sparse pattern (fig. A3
NVIDIA's SWS
Our SAS could utilize the same hardware by the novel structured sparse projection mecha- nism. With the same computational budget and on the same hardware, SAS realizes better accuracy if one can use extra memory for storing the weight.
S0LiUve Spars Virg!
Monocar AMgge Vnciiiaa
Calppl Actottionn
Figure A3: SWS matmul Sparse Tensor Core NVIDIA matmul mechanism in fig. 2
2020
Compare with our SAS
18


Reviews:
[{"review_id": "D0MBjEyKfw", "reviewer": "Reviewer_C5D8", "summary": "This paper introduces Structured Activation Sparsification (SAS), a method that enhances the accuracy of wide neural networks without the additional computational cost typically associated with network width. By implementing structured sparsity within activation maps—where a set number of non-zero values are maintained in consecutive activation. This allows for the simplification of wide matrix multiplications into narrow ones. Empirical results show that this method can improve accuracy (by up to 7% on CIFAR10) without increasing computational demands and outperforms similar sparsity approaches applied to network weights.", "questions": "I do not have other questions. Please refer to the weakness column for my comments.", "limitations": "", "rating": 6, "confidence": 3, "soundness": 3, "presentation": 3, "contribution": 3, "strengths": "The strengths of this work is listed below:\n1. This work introduces a novel method of structuring sparsity in activations, which appears to enable a reduction in computation without a corresponding drop in accuracy. The concept of using structured sparse projection (SAS) to maintain vectorization compatibility is particularly innovative.\n2. The structured sparsity allows for efficient matrix multiplication operations that maintain the number of multiplications at the level of the base dense/narrow network, highlighting efficiency in computation.\n3. The process for creating sparse activations through the structured sparse projection S is described as having negligible computational cost and wall-clock latency, indicating an efficient method that does not add significant complexity or processing time.\n4. A thorough evaluation is presented that demonstrates the SAS network's increased expressiveness compared to Static Weight Sparsity (SWS) networks, given the same computational budget, by using trajectory length analysis.\n5. The proposed SAS projection and its integration into the training process suggest a straightforward transformation of existing neural networks to increase their efficiency, which could be widely applicable across different network architectures and tasks.", "weaknesses": "The weakness of this work is listed below:\n1. As mentioned in Section 2.3, while the computational load in terms of FLOPS remains the same for a given level of activation sparsity \nM, the memory requirements for the Sparse Activation Sparsification (SAS) network increase linearly with M. This is in contrast to the Sparse Weight Storage (SWS) network, which maintains a constant storage requirement for weights at inference time regardless of M. This could become a significant limitation for devices with limited memory or when scaling to very large networks.\n2. In Section 3, the paper discusses the expressive power of SAS networks by comparing trajectory lengths in a specific constructed neural network with 2-dimensional input and output. However, this analysis might not generalize to all network architectures or datasets, which could limit the understanding of the practical implications of SAS's increased expressive power.\n3. The straightforward method used for computing the index I might not be the most effective approach, particularly when neighbor elements oscillate around zero (Section 6.1). An indexing strategy that is not learned end-to-end may limit the model's capacity to adapt to the data's complexity, potentially leaving some performance on the table.\n4. In general, the listings (code pieces) in the paper is informative enough. However, I would suggest the authors to replace the source code with some high-level pseudo code. This is more readable and more accessible to some readers.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "NVQqZtlmls", "reviewer": "Reviewer_W2qU", "summary": "This paper proposed SAS, a method to explore structured sparse activation in CNN. SAS describes the approach to generate structured sparse activation, software and hardware implementations. Numerical experiments on image classification validates the efficacy regarding accuracy aspect.", "questions": "- What is the training cost to yield SAS network compared to training as standard?", "limitations": "", "rating": 6, "confidence": 4, "soundness": 3, "presentation": 3, "contribution": 2, "strengths": "- The paper is written well, and technically sound.\n- The study problem is interesting and may deliver real impact onto DNN speedup.", "weaknesses": "- The terminology of structured sparsity is misleading. We typically refer the N:M sparsity as \"semi structured sparsity\" to distinguish it from standard structured sparsity including disjoint group sparsity, overlapping group sparsity and hierarchical sparsity. \n\n- The realistic benefits of structured sparse activation is not clear. Although the topic is interesting, I am not sure what is the actual speedup gain of such sparse activation that can deliver to the community. The paper seems equipping without numerical results regarding speedup as well.\n\n- The citation format is wrong. Please use \\citep{} rather than \\cite{} to cite references.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}, {"review_id": "VigHoxBPQ8", "reviewer": "Reviewer_N3vE", "summary": "This paper proposes a new concept of sparsity that maps the input activation values to a sparse representation, then exploits Nvidia Ampere sparsity for sparse computation by widening the weights. By doing this, wider weights make the network have stronger representation ability, and computation remains consistent because of the n:m sparsity operation. The authors conducted experiments on CIFAR and ImageNet datasets to verify the effectiveness of the proposed method.", "questions": "Please see the weakness part.", "limitations": "", "rating": 6, "confidence": 3, "soundness": 2, "presentation": 2, "contribution": 3, "strengths": "1. The presented idea is clever and novel. It can effectively enhance the representation ability of the network under the same amount of computation compared with the vanilla network. I think this brings a nice insight to the field.\n2. The author developed a general matmul library for the proposed SAS using Sparse Tensor Core. This gives the method practical value in the field, which is appreciated.", "weaknesses": "1. It is inaccurate for the author to state that the computation of SWS and SAS are the same. I can agree that the computation of SAS and the original dense network are the same, but the computation of SWS is obviously higher than that of SAS.\n2. Let's continue considering this point. Figure 3 raises a big question for me. The author doesn't express the dimensions of the weights corresponding to SWS and SAS. Although it intuitively looks like SAS has a clear acceleration in the graph, I think this is unreasonable because the author also says the number of mult count in SAS is the same as the base dense/narrow network. This greatly reduces my enthusiasm for this paper.\n3. The current presentation pf experiments is very scattered. I think the author should provide a comparison of the full network’s computation, inference speed, and accuracy to give readers a clearer comparison. For instance, when compressing ConvNeXt-b, although the accuracy of SAS is higher than SWS, a comparison of SAS's operation count and inference time also needs to be provided.", "quality": null, "clarity": null, "significance": null, "originality": null, "strengths_and_weaknesses": "", "was_revised": false, "final_justification": "", "rebuttal": "", "claims_and_evidence": "", "methods_and_evaluation_criteria": "", "theoretical_claims": "", "experimental_designs_or_analyses": "", "relation_to_broader_scientific_literature": "", "essential_references_not_discussed": "", "other_strengths_and_weaknesses": "", "other_comments_or_suggestions": "", "questions_for_authors": "", "summary_of_paper": "", "summary_of_recommendation": "", "technical_quality": "", "clarity_of_presentation": "", "potential_impact": "", "robotics_focus": "", "quality_of_limitations": "", "questions_for_rebuttal": "", "reasons_to_accept": "", "reasons_to_reject": "", "questions_to_authors": "", "extra_scores": {}}]

Author rebuttal:

Decision: Accept (poster)