A Parallel Algorithm for Enumerating Joint Weightof a Binary Linear Code in Network Coding
Shohei Ando, Fumihiko Ino, Toru Fujiwara, and Kenichi Hagihara
Graduate School of Information Science and Technology
Osaka University
1-5 Yamadaoka, Suita, Osaka 565–0871, Japan
Email: {s-ando, ino, fujiwara, hagihara}@ist.osaka-u.ac.jp

Abstract—In this paper, we present a parallel algorithm forenumerating joint weight of a binary linear (n, k) code, aiming
at accelerating assessment of its decoding error probability for
network coding. To reduce the number of pairs of codewords to
be investigated, our parallel algorithm reduces dimension k byfocusing on the all-one vector included in many practical codes.
We also employ a population count instruction to compute joint
weight of codewords with a less number of instructions. Ouralgorithm is implemented on a multi-core CPU system and an
NVIDIA GPU system using OpenMP and compute unified device
architecture (CUDA), respectively. We apply our implementation
to a subcode of a (127,22) BCH code to evaluate the impact ofacceleration.

I. I NTRODUCTION
Network coding [1] is a technique for improving transmission efficiency of multicast communication. This techniqueallows relay nodes to apply coding arithmetic to incoming
messages. Figure 1 shows an example of multicast communication using network coding over the butterfly network. In
this example, the source node s transmits two messages x andy to sink nodes r1 and r2 . On a typical network in Fig. 1(a),where relay nodes are prohibited to perform coding arithmetic,
r1 fails to receive y if v2 transmits x rather than y. Similarly,
r2 fails to receive x if v2 transmits y. In contrast, network
coding increases multicast efficiency by allowing relay nodev2 to transmit x⊕y to the next node v4 , as shown in Fig. 1(b).
The operator ⊕ here represents bitwise exclusive disjunction.
The sink node r1 then can extract y from its two incoming
messages x and x ⊕ y. Similarly, r2 can receive both x andy. Li et al. [2] presented that the max-flow bound from thesource node to each sink node can be achieved if relay nodes
use a linear transformation as such coding arithmetic.
In practice, an error-correcting code [3] must be applied
to flowing messages to achieve robust communication against
noise. The performance of an error-correcting code C can beassessed by performance metrics such as the decoding error
probability and error-correcting capability. An error here occurs if a transmitted message is decoded to a codeword v ∈ Cthat differs from the originally sent codeword u ∈ C (u ̸= v).
High-performance error correction can be achieved by not only
maximizing the error-correcting capability but also minimizing
the decoding error probability.
For a typical network, where relay nodes transmit incoming
messages without applying coding arithmetic, weight distribu-

s

x
v1

y
y

x

x
v3

v1

x
r1

y

x⊕ y

x

x

x⊕ y
r2

(a)

v3

v2

y
v4

y
y

x

v2

x

x

s

v4

r1

x⊕ y
r2

(b)

Fig. 1. An example of multicast communication using network coding overthe butterfly network. (a) Node r1 fails to receive y if node v2 transmits x
rather than y. (b) Network coding allows nodes r1 (r2 ) to receive both x andy, because y (x) can be extracted from x (y) and x ⊕ y.

tion of C is useful to evaluate the decoding error probability.
The weight distribution of C is denoted by an (n + 1)tuple (A0 , A1 , . . . , An ), where n is the length of C andAi (0 ≤ i ≤ n) is the number of codewords of Hamming
weight i in C [4]. This performance metrics is also useful toassess the performance of codes for network coding. However,a single error occurred on a network link can affect decoding
results of multiple messages. For example, the sink node r1 in
Fig. 1(b) can face with two incorrectly transmitted messages x
and x⊕y if an error occurs on the link between s and v1 . In thiscase, errors in incoming messages are dependent. Therefore,
the decoding error probability for m received codewords
cannot be obtained from weight distribution of C. Instead ofweight distribution, it requires m-joint weight distribution of
C [5]. In this work, we deal with the problem of the butterfly
network, so that we assume that m = 2 hereafter.
To the best of our knowledge, there is no formula that
directly gives joint weight distribution of a code C except
for Hamming code, simplex code, and the first order ReedMuller code [3]. Consequently, a joint weight enumerator isneeded to compute a joint weight histogram, which storesthe occurrence of joint weight for all 2-tuples (i.e., pairs)(u, v) (u, v ∈ C) of codewords in C. A histogram here
is an estimate of the probability distribution of a variable,
and consists of a sequence of bins, which store the frequencyof observations over categories (i.e., intervals of a variable).
Consider a binary linear (n, k) code [6] of length n anddimension k. The binary linear (n, k) code consists of 2kcodewords. Its joint weight histogram can be computed inO(22k n) time, because there are 22k pairs of codewords oflength n. This time complexity exponentially increases withk, so that the joint weight enumeration must be accelerated toassess codes with large dimension k.
In general, parallel-based solutions using multi-core CPUsor graphics processing units (GPUs) [7] are attractive methods
for achieving acceleration for compute- and memory-intensive
applications [8], [9], [10], [11]. Ando et al. [12] presented aparallel algorithm that enumerated joint weight on a multi-core
CPU and a GPU. They exploited data parallelism in enumeration by assigning different pairs of codewords to threads. Theiralgorithm employed an efficient mutual exclusion mechanism,
because multiple threads could simultaneously update thesame bin of the histogram. Furthermore, joint weight was
rapidly computed by a population count instruction. However,this parallel algorithm can be further accelerated by exploiting
theoretical properties on code structure.
In this paper, we propose a parallel algorithm for enumerating joint weight of a binary linear (n, k) code. We extend theprevious algorithm [12] by taking advantage of code structure
to reduce the number of pairs of codewords to be investigated.
We focus on the fact that many practical codes include theall-one vectors as a codeword. This assumption reduces thedimension k of the code for efficient enumeration. Similar
to [12], our algorithm employ a population count instruction
to rapidly compute joint weight of codewords. Our parallel
algorithm currently runs on a multi-core CPU system and acompute unified device architecture (CUDA) compatible GPU
system [13]. We assume that n ≤ 128 and the target machineis equipped with an NVIDIA Kepler GPU [14].
The following paper is organized as follows. Section II
introduces some related studies. Section III presents preliminaries on joint weight enumeration of a code. Section IV
describes our parallel algorithm for enumerating joint weightof a binary linear (n, k) code. Section V presents experimental
results. Finally, Section VI shows conclusion and future work.
II. R ELATED W ORK
Ando et al. [12] implemented a parallel algorithm thatenumerated joint weight of a binary linear (n, k) code on amulti-core CPU. Their CPU-based implementation exploited
multiple CPU cores by OpenMP [15], which achieved multithreading by simply adding compiler directives to the serial
code. Furthermore, single instruction multiple data (SIMD)
instructions, called Streaming SIMD Extensions (SSE) [16],were used to maximize the performance per CPU core byprocessing a 128-bit vector data simultaneously. They also presented a GPU-based implementation that processed thousandsof threads simultaneously. Their GPU-based implementation
generated pairs of codewords such that threads could accessdifferent memory addresses simultaneously. In contrast to
the architecture-specific optimization mentioned above, thepresent work focuses on code-oriented optimization that accelerate joint weight enumeration on arbitrary architectures.

Some theoretical results are useful to accelerate joint weight
enumeration of a binary linear (n, k) code. The MacWilliams
identity [3] is a famous theorem that relates the weight
enumerator of a binary linear (n, k) code to that of its dual
code, namely a binary linear (n, n − k) code. According tothis theorem, the time complexity of joint weight histogram
computation can be reduced from O(22k n) to O(22(n−k) n)
when k > n − k. Kido et al. [5] applied the MacWilliams
identity to an m-joint weight enumerator from a theoreticalpoint of view.
With respect to weight distribution, Desaki et al. [17]presented a weight enumeration algorithm that exploited code
structure called trellis diagram. Although this algorithm cannot
produce joint weight distribution of a code, their idea can beextended to joint weight enumeration algorithms to reduce theamount of work.
III. P RELIMINARIESLet u = (u1 u2 · · · un ) ∈ Fn and v = (v1 v2 · · · vn ) ∈ Fn
be vectors of length n, where F is a binary finite field andur , vr ∈ F (1 ≤ r ≤ n). Let fpq (u, v) also be the numberof r such that ur = p and vr = q, where p, q ∈ F. The joint
weight w(u, v) of a pair (u, v) of vectors then is given by a
4-tuplew(u, v) = (a, b, c, d),(1)where a = f11 (u, v), b = f10 (u, v), c = f01 (u, v), and d =f00 (u, v). For instance, we obtain (a, b, c, d) = (2, 2, 3, 1) foru = (11110000) and v = (11001110). Since d = n − a −b − c [3], we can omit the last element d from joint weight(a, b, c, d). Hereafter, we denote a joint weight by a 3-tuple(a, b, c) for simplicity.
Because a, b, and c are numbers enumerated from nelements, we have0 ≤ a, b, c ≤ n,
0 ≤ a + b + c ≤ n.

(2)
(3)

Joint weight enumeration outputs a sequence of numbers, eachcorresponding to the frequency of a tuple (a, b, c) that satisfies
Eqs. (2) and (3). Considering all combinations( with) repetitions,the number of possible tuples is given by n+3
4 .
Let C be a binary linear code. Let Ja,b,c be the numberof pairs (u, v) (u ∈ C, v ∈ C) of codewords that have joint
weight (a, b,(c). The
joint weight distribution of C is then)given by a n+3-tuple(J0,0,0 , J0,0,1 , . . . , Ja,b,c , . . . , Jn,0,0 )4such that a, b, and c satisfy Eqs. (2) and (3). This( distribution)
can be stored in a joint weight histogram with n+3
bins.
4
Algorithm 1 shows a brute-force parallel algorithm that generates joint weight histogram J from the length n and dimension k of code C, and a sequence W = (w0 , w1 , . . . , w2k −1 )of codewords in C. Parallelization can be easily achieved byassigning different pairs of codewords to threads. However, an
atomic instruction [13] is needed to compute the histogram
correctly, because multiple threads can simultaneously update the same bin at line 7. An alternative approach is toallow threads to have their own local histogram to prevent

Algorithm 1 Brute-force enumeration (J, n, k, W )Input: Length n, dimension k, and sequence W(w0 , w1 , . . . , w2k −1 ) of codewords
Output: Joint weight histogram J=

1: Initialize J;
2: for i ← 0 to 2k − 1 do parallelfor j ← 0 to 2k − 1 do parallela ← f11 (wi , wj );5:b ← f10 (wi , wj );6:c ← f01 (wi , wj );7:
Ja,b,c ← Ja,b,c + 1;
8:
end for
9: end for
3:
4:

1
2
3
4
5
6
7
8

int popcount(unsigned int x){
x=(x&0x55555555)+(x>>1&0x55555555);
x=(x&0x33333333)+(x>>2&0x33333333);
x=(x&0x0f0f0f0f)+(x>>4&0x0f0f0f0f);x=(x&0x00ff00ff)+(x>>8&0x00ff00ff);
return (x&0x0000ffff)+(x>>16&0x0000ffff);}

Fig. 2.

simultaneous access to the same memory address. Although
this approach avoids atomic instructions, a post-processing
stage is needed to merge local histograms into a single global
histogram. More details on this hierarchical organization arepresented in [12].
For a multi-core CPU system, the first loop at line 2 canbe parallelized using multiple threads by adding an OpenMP
directive [15] such as #pragma omp parallel for. On
the other hand, the nested loop structure from lines 2 to 9
can be replaced with a kernel function call for GPU-based
acceleration. The kernel function implements the loop bodyfrom lines 4 to 7 to enumerate joint weight in parallel.
A. Joint weight computation
Joint weight w(u, v) = (a, b, c) of a pair (u, v) of codewords can be given by
a = popcount(u ∧ v),
b = popcount(u) − a,

(4)
(5)

c = popcount(v) − a,(6)where ∧ is bitwise logical conjunction and popcount(u) is afunction that counts the number of bits set to 1 in the given
vector u. Figure 2 shows a naive implementation [18] of the
popcount function. This implementation processes a vector ofl = 32 bits in log l steps, which require 20 instructions containing five additions, five shifts, and ten logical conjunctions.
Instead of this implementation, the previous algorithm [12]employed a single vector instruction to process a vector of 64
bits. For multi-core CPUs and GPUs, the POPCNT instructionof SSE 4.2 [16] and the popc instruction of CUDA [13],respectively, were used to implement the popcount function.
These vector instructions require 64-bit data as their operand,so that ⌈n/64⌉ instructions are needed to compute joint weightof a pair of codeword of length n.
B. Reduced enumeration using symmetry
Let w(u, v) = (a, b, c) be the joint weight of pair (u, v) ofcodewords. The joint weight of the permutated pair (v, u) thenis given by w(v, u) = (a, c, b) [3]. This symmetric relation

Basic implementation [18] of the popcount function for 32-bit data.

Algorithm 2 Symmetric enumeration (J, n, k, W )Input: Length n, dimension k, and sequence W(w0 , w1 , . . . , w2k −1 ) of codewords
Output: Joint weight histogram J
1: Initialize J, D, and I;
2: for i ← 0 to 2k − 1 do parallel3:
for j ← i to 2k − 1 do parallel4:a ← f11 (wi , wj );5:b ← f10 (wi , wj );6:c ← f01 (wi , wj );7:
if i = j then8:
Da,b,c ← Da,b,c + 1;
9:
else
10:
Ia,b,c ← Ia,b,c + 1;
11:
end if
12:
end for
13: end for
14: for a ← 0 to n do
15:
for b ← 0 to n − a do
16:
for c ← 0 to n − a − b do
17:Ja,b,c ← Ia,b,c + Ia,c,b + Da,b,c ;18:
end for19:
end for
20: end for=indicates that joint weight distribution can be obtained fromapproximately half of pairs of codewords.
Consider a binary linear (n, k) code C that contains 2k
codewords w0 , w1 , . . . , w2k −1 . Let Ia,b,c be the number ofpairs of codewords such that (wi , wj ) has joint weight ofw(wi , wj ) = (a, b, c), where 0 ≤ i < j < 2k . Let Da,b,c also
be the number of pairs of codewords such that (wi , wi ) has
joint weight of w(wi , wi ) = (a, b, c), where 0 ≤ i < 2k . We
then have the following relation:
Ja,b,c = Ia,b,c + Ia,c,b + Da,b,c .

(7)

I and D can be obtained from 2k (2k − 1)/2 and 2k pairsof codewords, respectively. Therefore, this symmetric relation
reduces the number of pairs of codewords to be investigatedfrom 22k to 2k (2k + 1)/2 ≈ 22k−1 .
Algorithm 2 shows a joint weight enumeration algorithm
that exploits this symmetry. The nested loop from lines 2 to
13 computes a joint weight histogram for approximately halfof pairs of codewords. The entire histogram is then serially

Algorithm 3 Conflict tolerant enumeration (J, n, k, W )Input: Length n, dimension k, and sequence W(w0 , w1 , . . . , w2k −1 ) of codewords
Output: Joint weight histogram J=

1: Sort W in ascending order in terms of Hamming weight;
2: W ′ ← ∅;3: while (W ̸= ∅) do4:
for i ← 0 to n do

if (∃u ∈ W such that popcount(u) = i) thenDelete u from W ;7:
Add u to W ′ ;8:
end if
9:end for
10: end while
11: Call symmetric enumeration (J, n, k, W ′ );5:
6:

computed from lines 14 to 20 using Eq. (7).
C. Conflict tolerant enumeration on GPU
As compared with the CPU, the GPU is a highly-threaded
architecture capable of running thousands of threads simultaneously. Therefore, the overhead of atomic instructions can bea performance bottleneck in GPU-based enumeration. To deal
with this issue, the previous GPU-based implementation [12]reduced write conflicts by scanning pairs of codewords suchthat different threads updated different bins. The symmetry
mentioned in Section III-B was used to realize such conflict
tolerant enumeration.
Suppose that pairs (u, ∗) and (u′ , ∗) of codewords are assigned to threads #1 and #2, respectively, where ∗ is an arbitral
codeword to be investigated. Suppose that popcount(u) = a+b
and popcount(u′ ) = a′ + b′ . We have a ̸= a′ or b ̸= b′ if
popcount(u) ̸= popcount(u′ ). In this case, threads #1 and #2
update a different bin, and thus, write conflicts do not occur
between them. Therefore, codewords to be investigated should
be classified into groups in terms of popcount value (i.e.,
Hamming weight). Threads #1 and #2 then are responsible forpairs (u, ∗) and (u′ , ∗) of codewords such that popcount(u) ̸=popcount(u′ ).
A preprocessing stage is required to realize this assignment.
That is, codewords should be sorted in ascending order interms of Hamming weight. This preprocessing stage canbe processed in O(2k n) time, which is much smaller thanO(22k n), or the time complexity of joint weight distribution
computation. Consequently, this sorting operation is processedon a CPU. After this preprocessing stage, joint weight distribution is computed using Algorithm 2.
Algorithm 3 shows a pseudocode of the previous algorithm
for GPU-based enumeration [12]. The preprocessing stage
from lines 1 to 10 produces a sequence W ′ of sorted codewords, which is then given as an input to Algorithm 2.
IV. P ROPOSED J OINT W EIGHT E NUMERATION
Our parallel joint weight enumeration algorithm consists offour acceleration techniques: (1) dimension reduction using the

all-one vector, (2) joint weight computation with a population
count instruction (Section III-A), (3) reduced enumeration
using symmetry (Section III-B), and (4) conflict tolerant
enumeration on the GPU (Section III-C).
A. Dimension reduction using all-one vector
Many practical codes such as Bose-ChaudhuriHocquenghem (BCH) codes [19], [20], Hamming codes,
and Reed-Muller codes include the all-one vector 1 as acodeword. Given such code C, we haveu + 1 = u ∈ C, for all u ∈ C.

(8)This implies that dimension k can be reduced for enumerating
joint weight efficiently. That is, w(u + 1, v), w(u, v + 1), andw(u + 1, v + 1) can be obtained from w(u, v) = (a, b, c, d)
as follows:w(u + 1, v) = (c, d, a, b),w(u, v + 1) = (b, a, d, c),w(u + 1, v + 1) = (d, c, b, a).

(9)
(10)
(11)

In other words, given codewords u and v, there is no need togenerate codewords u + 1 and v + 1 to compute their joint
weight. Therefore, the number of pairs of codewords can be
reduced to quarter.
B. Codeword generation with dimension reduction
Before using the dimension reduction technique mentioned
above, we have to confirm whether the target code C includes
the all-one vector 1 or not. A straightforward scheme toperform this confirmation is to use a parity-check matrix H of
C: a vector x is a codeword of C if H·xT = 0, where xT is thetranspose of x. Otherwise, the dimension reduction techniquecannot be used for enumeration. Although this scheme detects
the all-one vector in C, it is not clear which codeword mustbe generated for enumeration.
Our alternative scheme is as follows. Let G be the canonicalgenerator matrix [3] of a binary linear (n, k) code C. G canbe represented as G = [ Ik |P ], where Ik is the k × k identity
matrix and P is a k × (n − k) matrix.
1) Replace∑ the bottom row Gk of G with summation ofrows, 1≤i≤k Gi , where Gi represents the i-th row ofG.
2) Check the bottom row Gk to detect the all-one vector.
a) If Gk = 1T , code C includes the all-one vector 1. We then eliminate Gk from G, because
codewords u + 1 and v + 1 are not needed tocompute their joint weight, as mentioned in Section
IV-A. This elimination reduces the dimension ofC, because the generator matrix is then given by
G = [ Ik−1 |P ′ ], where P ′ is a (k − 1) × (n − k)matrix.
b) Otherwise, code C does not include the all-one
vector 1.
3) Use G to generate the codewords of C.

Item
Machine #1
Machine #2# of CPU sockets
1
2
CPUCore i7 3770K
Xeon E5-2680v2# of cores per socket
410Frequency
3.5 GHz
2.8 GHz
Main memory capacity
16 GB
512 GB
Peak memory bandwidth
25.6 GB/s
51.2 GB/sGPU
GTX 680
Tesla K40
# of cores
1536
2880
Core clock
1006 MHz745 MHz
VRAM capacity
2 GB
12 GB
Peak memory bandwidth
192 GB/s
288 GB/s
PCI Express bus
3.0 ×16
3.0 ×16OS
Ubuntu 12.04 64-bit Ubuntu 13.10 64-bit
GCC 4.6.3
GCC 4.8.1
C++ compiler
Graphics driver331.62CUDA
6.0-arch sm_30 -O3 -arch sm_35 -O3Compile option

Execution time (s)TABLE IS PECIFICATION OF EXPERIMENTAL MACHINES .

5,000
4,500
4,000
3,500
3,000
2,500
2,000
1,500
1,000
500
0

Previous
This work

Core i7

Xeon E5

GTX 680

K40

Fig. 3. Execution times of the proposed algorithm and the previous algorithm[12].

To evaluate the performance of our parallel algorithm, we
compared our implementation with the previous implementation [12] in terms of execution time. Table I shows the
specification of our experimental machines. Our machines
had a 4-core Core i7 CPU and two 10-core Xeon E5 CPUs,
respectively. In addition, these machines were equipped with a
GTX 680 GPU and a Tesla K40 GPU, respectively. Both GPUs
were based on the Kepler architecture [14]. The error checkand correct (ECC) capability of the K40 card was turned off
during measurement. The algorithms were implementated using OpenMP to realize multithreaded enumeration. We created
two threads per CPU core to take advantage of hyperthreading
technology.
We used a subcode of the (127,22) binary BCH code [19],[20] of dimension t (≤ k), where 11 ≤ t ≤ 22. Consequently,
a joint weight histogram was computed for 2t codewords. The
bin size of the global joint weight histogram was set to 8 bytes,
because the maximum value of a bin could reach 22t , where
11 ≤ t ≤ 22. On the other hand, the bin size of local joint
histograms was set to 2 bytes and 4 bytes for the CPU- andGPU-based implementations, respectively [12].

Because the K40 card had a higher memory bandwidth thanthe GTX 680 card, we expected that the former would outperform the latter for this memory-intensive application. However,
we found that the GTX 680 card was 1.36 times faster thanthe K40 card for this application. Because this speedup ratio is
close to the clock speed ratio (1006/745 = 1.35), we think that
this performance behavior is due to atomic instructions. The
throughput of atomic instructions on the Kepler architecture
varies according to the memory address to be accessed. For
example, the worst throughput of 1 instruction per clock is
obtained when threads access the same address simultaneously[21]. In contrast, this throughput increases to 8 instructions
per streaming multiprocessor [13] per clock when threads inthe same warp access different addresses in the same cacheline. Consequently, the worst throughput rather than the best
throughput determines the performance of our GPU-based
implementation, though write conflicts are reduced by sorting
codewords.
Our algorithm requires a preprocessing stage on the CPU.
Both the CPU- and GPU-based implementations generate
codewords before computing the joint weight histogram. The
GPU-based implementation then sorts codewords and transfers
them to the GPU. We found that the preprocessing overhead
was negligible against the entire execution time. For example,
when t = 22, the preprocessing time and the transfer time
were approximately 0.009% and 0.002% of the execution
time on the GTX 680 card, respectively. The performanceof our algorithm is dominated by joint weight distribution
computation.

A. Performance Comparison

B. Efficiency Analysis

Figure 3 shows the execution times for n = 127 and t = 22.
The execution times of GPU-based implementations include
the transfer time needed to copy data between the CPU andGPU. On all CPUs and GPUs, our algorithm reduced theexecution time to quarter. These timing results are the same aswhat expected in Section IV-A. The highest performance wasobtained on the Xeon processors, which reduced the execution
time from 1,111 to 275 seconds. The speedup over the previous
algorithm reached a factor of 4.03, which was slightly higher
than a factor of 4.

We next analyzed the efficiency of the implementations interms of memory throughput, because access to histogram
bins determined their performance. The effective memory
throughput is given by B = AM/T , where A, M , and T
are the number of pairs of codewords, the amount of memory
reads/writes per pair, and the execution time, respectively. Notethat the execution time T here includes the preprocessing
time and the transfer time mentioned above. For each pair ofcodewords of length n, our implementation updates an 8-bytebin. Therefore, we have M = 2 ⌈n/8⌉+8 in byte. Considering

Similar to the sorting operation presented in Section III-C,
we decided to serially process this preprocessing stage on aCPU.
V. E XPERIMENTAL R ESULTS

20Core i7300

16

Xeon E5250

GTX 680

200

K40

Speedup

Memory throughput (GB/s)

350

150

12
8

100

4

50

Core i7

Xeon E5

Linear

28

36

0

0
11

12

13

14

15

16 17 18
Dimension t

19

20

21

1

22

Fig. 4. Effective memory throughputs measured during parallel joint weight
enumeration.

combinations of two codewords( k )to be selected from the total
2k codewords, we have A = 22 .
Figure 4 shows the effective memory throughputs of the implementations. When t = 22, The effective memory throughputs of the GTX 680 and K40 reached 191 GB/s and 141
GB/s, respectively. These results were equivalent to 99.6%
and 49.0% of the peak memory bandwidth, respectively. The
former is close to 100%, demonstrating the effectiveness ofhierarchical histogram organization mentioned in Section III.
That is, local histograms are small enough to fit into the
GPU cache called shared memory [13]. Therefore, the memory
bandwidth was efficiently saved using the shared memory,increasing the effective memory throughput close to the peakmemory bandwidth.
With respect to CPU-based results, we found that the effective memory throughputs were higher than the peak memory
bandwidth. Similar to GPU-based results, this behavior canbe explained by cache hits. Actually, a local joint weight
histogram for n ≤ 128 can be stored in a memory regionof approximately 17 KB [12], which is smaller than thecapacity of L1 cache (32 KB). Consequently, the instruction
issue rate determines the performance of our CPU-based
implementation. According to this analysis, the two Xeon
processors are 4 times faster than the single Core processor,
because the formers have five times more physical cores butwith 20% slower clock rate than the latter. Actually, the Xeon
processors achieved approximately 4.3 times higher memory
throughput than the Core processor when t = 22.
When t ≤ 16, the Xeon processors failed to outperform theCore processor. This is due to the memory allocation overhead.
For example, memory allocation on this big memory machine
took 0.4 seconds, whereas the execution time at t = 16 was
0.76 seconds.
Finally, Fig. 5 shows the speedup over a single-core version
on the Core and Xeon processors, which have 4 and 20
physical cores, respectively. The performance of our implementation linearly increased with the number of threads.
This performance behavior also explains why the memory
throughputs are higher than the peak memory bandwidth.
Because local histograms were entirely stored in the L1 cache,the performance was mainly dominated by the instruction issue

Fig. 5.

48

12

16 20 24
# of threads

32

40

Speedups over a single-core version on Core and Xeon processors.

rate. Therefore, the performance increased with the number ofphysical cores to be exploited for enumeration. Owing to thehyperthreading technology, the performance slightly increased
after assigning the second threads to CPU cores.
VI. C ONCLUSION
In this paper, we presented a parallel algorithm for enumerating joint weight of a binary linear (n, k) code. Our algorithm
reduces the number of pairs of codewords to be investigated.
To realize this, we reduce the dimension k of the code by focusing on the all-one vector, which is included in typical errorcorrecting codes. Our algorithm also employ a population
count instruction to reduce the number of instructions needed
to compute joint weight. We also sort codewords in terms ofHamming weight to realize conflict tolerant enumeration.
Our experimental results showed that the dimension reduction reduced the execution time to quarter on multi-core CPUs
and a GPU. We also found that the performance of our GPUbased implementation was dominated by the core clock speedof the GPU. Similarly, our CPU-based implementation had aperformance bottleneck in the instruction issue rate.
Future work includes further exploitation of code structuresuch as trellis diagram [17]. The MacWilliams identity is also
useful to accelerate enumeration for codes of large dimension.
ACKNOWLEDGMENT
This study was supported in part by the Japan Society
for the Promotion of Science KAKENHI Grant Number
24560458 and the Japan Science and Technology Agency
CREST program, “An Evolutionary Approach to Construction of a Software Development Environment for MassivelyParallel Computing Systems.”
R EFERENCES[1] R. Ahlswede, N. Cai, S.-Y. R. Li, and R. W. Yeung, “Network information flow,” IEEE Trans. Information Theory, vol. 46, no. 4, pp. 1204–
1216, Jul. 2000.
[2] S.-Y. R. Li, R. W. Yeung, and N. Cai, “Linear network coding,” IEEE
Trans. Information Theory, vol. 49, no. 2, pp. 371–381, Feb. 2003.
[3] F. J. MacWilliams and N. J. A. Sloane, The Theory of Error-Correcting
Codes. North-Holland, 1977.
[4] M. Mohri, Y. Honda, and M. Morii, “A method for computing the local
distance profile of binary cyclic codes,” IEICE Trans. Fundamentals(Japanese Edition), vol. J86-A, no. 1, pp. 60–74, Jan. 2003.

[5] Y. Kido and T. Fujiwara, “MacWilliams identity for joint weight
enumerator to evaluate decoding error probability of linear block codein network coding,” in Proc. 34th Symp. Information Theory and its
Applications (SITA’11), 3.4.1, Nov. 2011, pp. 184–189, (In Japanese).
[6] S. Lin and D. J. Costello, Error Control Coding: Fundamentals andApplications, 2nd ed. Prentice Hall, 2004.
[7] D. Luebke and G. Humphreys, “How GPUs work,” Computer, vol. 40,
no. 2, pp. 96–100, Feb. 2007.
[8] M. Garland, S. L. Grand, J. Nickolls, J. Anderson, J. Hardwick,
S. Morton, E. Phillips, Y. Zhang, and V. Volkov, “Parallel computing
experiences with CUDA,” IEEE Micro, vol. 28, no. 4, pp. 13–27, Jul.
2008.
[9] K. Ikeda, F. Ino, and K. Hagihara, “Efficient acceleration of mutual
information computation for nonrigid registration using CUDA,” IEEE
J. Biomedical and Health Informatics, vol. 18, no. 3, pp. 956–968, May
2014.
[10] F. Ino, Y. Munekawa, and K. Hagihara, “Sequence homology searchusing fine grained cycle sharing of idle GPUs,” IEEE Trans. Paralleland Distributed Systems, vol. 23, no. 4, pp. 751–759, Apr. 2012.
[11] Y. Okitsu, F. Ino, and K. Hagihara, “High-performance cone beamreconstruction using CUDA compatible GPUs,” Parallel Computing,vol. 36, no. 2/3, pp. 129–141, Feb. 2010.
[12] S. Ando, F. Ino, T. Fujiwara, and K. Hagihara, “A parallel method
for accelerating joint weight distribution computation,” IEICE Trans.
Information and Systems (Japanese Edition), vol. J97-D, no. 9, pp. xx–
xx, Sep. 2014, (In Japanese).
[13] NVIDIA Corporation, “CUDA C Programming Guide Version 6.0,” Feb.
2014, http://docs.nvidia.com/cuda/pdf/CUDA C Programming Guide.
pdf.
[14] ——, “NVIDIA’s Next Generation CUDA Compute Architecture: Kepler GK110,” May 2012, http://www.nvidia.com/content/PDF/kepler/
NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf.
[15] R. Chandra, L. Dagum, D. Kohr, D. Maydan, J. McDonald, and
R. Menon, Parallel Programming in OpenMP. San Mateo, CA: Morgan
Kaufmann, Oct. 2000.
[16] Intel Corporation, “Intel SSE4 Programming Reference,” Jul. 2007, http://software.intel.com/sites/default/files/m/8/b/8/D9156103.pdf.
[17] Y. Desaki, T. Fujiwara, and T. Kasami, “A method for computing theweight distribution of a block code by using its trellis diagram,” IEICE
Trans. Fundamentals, vol. E77-A, no. 8, pp. 1230–1237, Aug. 1994.
[18] H. S. Warren, Hacker’s Delight, 2nd ed. Addision-Wesley Professional,
2012.
[19] R. C. Bose and D. K. Ray-Chaudhuri, “On a class of error correcting
binary group codes,” Information and Control, vol. 3, no. 1, pp. 68–79,
Mar. 1960.
[20] A. Hocquenghem, “Codes correcteurs d’erreurs,” Chiffres, vol. 29, no. 2,
pp. 147–156, Apr. 1959, (In French).
[21] L. Nyland and S. Johns, “Understanding and using atomic
memory operations,” in 4th GPU Technology Conf. (GTC’13),
Mar. 2013, http://on-demand.gputechconf.com/gtc/2013/presentations/
S3101-Atomic-Memory-Operations.pdf.