Data-Level Parallelism in Vector, SIMD, and GPU Architectures

4.1 Introduction

P282

three advantage of SIMD compared to MIMD

  1. is not only the matrix-oriented computations of scientific computing but also the media-oriented image and sound processing and machine learning algorithms
  2. MIMD architecture needs to fetch one instruction per data operation, single instruction multiple data (SIMD) is potentially more energy-efficient since a single instruction can launch many data operations
  3. biggest advantage of SIMD versus MIMD is that the programmer continues to think sequentially yet achieves parallel speedup by having parallel data operations.

So where is the description of MIMD in this textbook

Property vector architectures multimedia SIMD instruction set extensions graphics processing units (GPUs)
Advantage easier to understand and to compile to than other SIMD variations    
Disadvantage but they were considered too expensive for microprocessors until recently    

This environment has a system processor and system memory in addition to the GPU and its graphics memory. In fact, to recognize those distinctions, the GPU community refers to this type of architecture as heterogeneous.

P283

vector architectures are supersets of themultimedia SIMD instructions, including a better model for compilation

4.2 Vector Architecture

In practice, developers can express many of the programs that ran well on complex out-oforder designs more efficiently as data-level parallelism in the form of vector instructions

P284

The vector and scalar registers have a significant number of read and write ports to allow multiple simultaneous vector operations.

P285

primary components of the instruction set architecture of RV64V are the following:

  1. vector register

    these ports will allow a high degree of overlap among vector operations to different vector registers, what does this mean ?

  2. Vector functional units
  3. Vector load/store unit
  4. A set of scalar registers

An innovation of RV64V is to associate a data type and data size with each vector register, rather than the normal approach of the instruction supplying that information

P286

P287

Reasons for dynamic register typing:

  1. many instructions are required for a conventional vector architecture that supports such variety. 2.
  2. dynamic typing also lets programs disable unused vector registers

    As a consequence, enabled vector registers are allocated all the vector memory as long vectors. For example, assume we have 1024 bytes of vector memory, if 4 vector registers are enabled and they are type 64–bit floats, the processor would give each vector register 256 bytes or 256/8¼32 elements. This valiue is called the maximum vector length (mvl), which is set by the processor and cannot be changed by sofware

  3. A pleasant side effect of dynamic register typing is that the program can configure vector registers as disabled when they are not being used, so there is no need to save and restore them on a context switch.

    Our implementation of RV64V increases state a factor of 3: from 2 * 32 * 8 = 512 bytes to 2 * 32 * 1024 = 1536 bytes. !! i am good at arithmetic

  4. third benefit of dynamic register typing is that conversions between different size operands can be implicit depending on the configuration of the registers rather than as additional explicit conversion instructions

so what is vector register, is a vector of register or a just normal register whose size is 64bit

The vector-length register vl is used when the natural vector length is not equal to mvl, the vector-type register vctype records register types, and the predicate registers pi are used when loops involve IF statements.

P288

vsetdcfg 4*FP64 # Enable 4 DP FP vregs
fld f0,a # Load scalar a
vld v0,x5 # Load vector X
vmul v1,v0,f0 # Vector-scalar mult
vld v2,x6 # Load vector Y
vadd v3,v1,v2 # Vector-vector add
vst v3,x6 # Store the sum
vdisable # Disable vector regs

why vsetdcfg is 4 * PF64 instead of 32 * PF64

P289

When the compiler produces vector instructions for such a sequence, and the resulting code spends much of its time running in vector mode, the code is said to be vectorized or vectorizable. Loops can be vectorized when they do not have dependences between iterations of a loop, which are called loop-carried dependences

in the following two snippet, I still can not understand the vsetdcfg

Vector architects call forwarding of element dependent operations chaining, in that the dependent operations are “chained” together

P290

make a lot of assumption about how to calculate execution time and introduce some conception

The execution time of a sequence of vector operations primarily depends on three factors:

  1. the length of the operand vectors,
  2. structural hazards among the operations, and
  3. the data dependences.

To simplify the discussion of vector execution and vector performance, we use the notion of a convoy, which is the set of vector instructions that could potentially execute together

P291

Early implementations of chaining worked just like forwarding in scalar pipelining, but this restricted the timing of the source and destination instructions in the chain

how to understand the timing

Recent implementations use flexible chaining, which allows a vector instruction to chain to essentially any other active vector instruction, assuming that we don’t generate a structural hazard

any differences between early implement, we can implement both of them

To turn convoys into execution time, we need a metric to estimate the length of a convoy. It is called a chime, which is simply the unit of time taken to execute one convoy.

Thus a vector sequence that consists of m convoys executes in m chimes; for a vector length of n, for our simple RV64V implementation, this is approximately m * n clock cycles

vector !, vector !, so what is meaning of vector if exection of time is proportional to vection length

One source of overhead ignored inmeasuring chimes is any limitation on initiating multiple vector instructions in a single clock cycle

I can’t translate this into chinese, what do you mean by any limitation on initiating ?

Example

The first convoy starts with the first vld instruction. The vmul is dependent on the first vld, but chaining allows it to be in the same convoy

P292

The most important source of overhead ignored by the chime model is vector start-up time, which is the latency in clock cycles until the pipeline is full. The start-up time is principally determined by the pipelining latency of the vector functional unit

Q A
How can a vector processor execute a single vector faster than one element per clock cycle? Multiple elements per clock cycle improve performance.
How does a vector processor handle programs where the vector lengths are not the same as the maximum vector length (mvl)? Because most application vectors don’t match the architecture vector length, we need an efficient solution to this common case.
What happens when there is an IF statement inside the code to be vectorized? More code can vectorize if we can efficiently handle conditional statements.
What does a vector processor need from the memory system ? Without sufficient memory bandwidth, vector execution can be futile.
How does a vector processor handle multiple dimensional matrices? This popular data structure must vectorize for vector architectures to do well.
How does a vector processor handle sparse matrices? This popular data structure must vectorize also.
How do you program a vector computer? Architectural innovations that are a mismatch to programming languages and their compilers may not get widespread use

P293

Multiple Lanes: Beyond One Element per Clock Cycle

P294

P295

P296

Vector-Length Registers: Handling Loops Not Equal to 32

RISC-V has a better solution than a separate loop for strip mining

P297

Predicate Registers: Handling IF Statements in Vector Loops The presence of conditionals (IF statements) inside loops and the use of sparse matrices are two main reasons for lower levels of vectorization

Compiler writers use the term IF-conversion to transform an IF statement into a straight-line code sequence using conditional execution.

one difference between vector processors and GPUs is the way they handle conditional statements

P298

To maintain an initiation rate of one word fetched or stored per clock cycle, the memory system must be capable of producing or accepting this much data. Spreading accesses across multiple independent memory banks usually delivers the desired rate

P299

Memory Banks: Supplying Bandwidth for Vector Load/Store Units Most vector processors use memory banks, which allow several independent accesses rather than simple memory interleaving for three reasons:

  1. Many vector computers support many loads or stores per clock cycle, and the memory bank cycle time is usually several times larger than the processor cycle time. To support simultaneous accesses from multiple loads or stores, the memory system needs multiple banks and needs to be able to control the addresses to the banks independently.
  2. Most vector processors support the ability to load or store data words that are not sequential. In such cases, independent bank addressing, rather than interleaving, is required.
  3. Most vector computers support multiple processors sharing the same memory system, so each processor will be generating its own separate stream of addresses.

P300

Stride: Handling Multidimensional Arrays in Vector Architectures This distance separating elements to be gathered into a single vector register is called the stride.

Caches inherently deal with unit-stride data; increasing block size can help reduce miss rates for large scientific datasets with unit stride, but increasing block size can even have a negative effect for data that are accessed with nonunit strides.

Supporting strides greater than one complicates the memory system. Once we introduce nonunit strides, it becomes possible to request accesses from the same bank frequently.

!!! can not understand the formula at page 301, maybe a typo

P301

Gather-Scatter: Handling Sparse Matrices in Vector Architectures

P302

P303

Programming Vector Architectures It is this dialogue between the compiler and the programmer, with each side giving hints to the other on how to improve performance, that simplifies programming of vector computers

4.3 SIMD Instruction Set Extensions for Multimedia

P304

SIMD Multimedia Extensions started with the simple observation that many media applications operate on narrower data types than the 32-bit processors were optimized for

Unlike vector machines with large register files SIMD instructions tend to specify fewer operands and thus use much smaller register files

In contrast to vector architectures, which offer an elegant instruction set that is intended to be the target of a vectorizing compiler, SIMD extensions have three major omissions

  1. no vector length register
  2. no strided or gather/scatter data transfer instructions
  3. no mask register

P305

P306

SIMD Multimedia Extensions started with the simple observation that many media Given these weaknesses, why are multimedia SIMD extensions so popular?

The overarching issue, however, is that due the overiding importance of backwards binary compatability, once an architecture gets started on the SIMD path it’s very hard to get off it.

To get an idea about what multimedia instructions look like, assume we added a 256-bit SIMD multimedia instruction extension to RISC-V, tentatively called RVP for “packed.”

Oh, great, RISC-V doesn’t have SIMD Instruction

P307

Programming Multimedia SIMD Architectures

P308

The Roofline Visual Performance Model

skip, because it made me feel bad

Graphics Processing Units

P310

Thus the design of GPUs may make more sense when architects ask, given the hardware invested to do graphics well, how can we supplement it to improve the performance of a wider range of applications?

P311

The challenge for the GPU programmer is

  1. not simply getting good performance on the GPU,
  2. but also in coordinating the scheduling of computation on the system processor and the GPU and the transfer of data between system memory and GPU memory.
  3. Moreover, as we see will see later in this section, GPUs have virtually every type of parallelism that can be captured by the programming environment: multithreading, MIMD, SIMD, and even instruction-level

NVIDIA decided that the unifying theme of all these forms of parallelism is the CUDA Thread.

Using this lowest level of parallelism as the programming primitive, the compiler and the hardware can gang thousands of CUDA Threads together to utilize the various styles of parallelism within a GPU: multithreading, MIMD, SIMD, and instruction-level parallelism.

Therefore NVIDIA classifies the CUDA programming model as single instruction, multiple thread (SIMT).

For reasons we will soon see, these threads are blocked together and executed in groups of threads, called a Thread Block. We call the hardware that executes a whole block of threads a multithreaded SIMD Processor

We need just a few details before we can give an example of a CUDA program:

  1. To distinguish between functions for the GPU (device) and functions for the system processor (host), CUDA uses __device__ or __global__ for the former and __host__ for the latter.
  2. CUDA variables declared with __device__ are allocated to the GPU Memory (see below), which is accessible by all multithreaded SIMD Processors.
  3. The extended function call syntax for the function name that runs on the GPU is 
    name < <<dimGrid, dimBlock>> > (… parameter list…)

    where dimGrid and dimBlock specify the dimensions of the code (in Thread Blocks) and the dimensions of a block (in threads).

  4. In addition to the identifier for blocks (blockIdx) and the identifier for each thread in a block (threadIdx), CUDA provides a keyword for the number of threads per block (blockDim), which comes from the dimBlock parameter in the preceding bullet.

P312

Different Thread Blocks cannot communicate directly, although they can coordinate using atomic memory operations in global memory.

As we will soon see, many GPU hardware concepts are not obvious in CUDA. Writing efficient GPU code requires that programmers think in terms of SIMD operations, even though the CUDA programming model looks like MIMD

Like many parallel systems, a compromise between productivity and performance is for CUDA to include intrinsics to give programmers explicit control over the hardware

P313

NVIDIA GPU Computational Structures One obstacle to understanding GPUs has been the jargon, with some terms even having misleading names.

ometimes, GPUs implement certain features in hardware that vector processors would implement in software. This difference is because vector processors have a scalar processor that can execute a software function

Unlike most vector architectures, GPUs also rely on multithreading within a single multithreaded SIMD Processor to hide memory latency

so, how vector architecture hide memory lantency ?

本站所有文章转发 CSDN 将按侵权追究法律责任，其它情况随意。