2016年12月简讯

这个月终于结束了将近三个月的失业状态,开始了新的工作。

工作方面:
开始接触和学习CUDAGPU编程,此外,又把Python捡了起来,写了一个稍微有点规模的小程序。

业余时间(技术):
写了一篇Unix之殇的用心之作。

业余时间(读书):
读了一本介绍拜仁慕尼黑俱乐部的画报。

生活方面:
参加了两次朋友家举办的party,结识了一些新朋友,也从他们那里了解到很多逸闻趣事。

CUDA编程笔记(15)——MEMORY ACCESS PATTERNS

这篇笔记摘自Professional CUDA C Programming

Global memory loads/stores are staged through caches, as shown in Figure 4-6. Global memory is a logical memory space that you can access from your kernel. All application data initially resides in DRAM, the physical device memory. Kernel memory requests are typically served between the device DRAM and SM on-chip memory using either 128-byte or 32-byte memory transactions.

All accesses to global memory go through the L2 cache. Many accesses also pass through the L1 cache, depending on the type of access and your GPU’s architecture. If both L1 and L2 caches are used, a memory access is serviced by a 128-byte memory transaction. If only the L2 cache is used, a memory access is serviced by a 32-byte memory transaction. On architectures that allow the L1 cache to be used for global memory caching, the L1 cache can be explicitly enabled or disabled at compile time.

An L1 cache line is 128 bytes, and it maps to a 128-byte aligned segment in device memory. If each thread in a warp requests one 4-byte value, that results in 128 bytes of data per request, which maps perfectly to the cache line size and device memory segment size.

There are two characteristics of device memory accesses that you should strive for when optimizing your application:
➤ Aligned memory accesses
➤ Coalesced memory

capture

Aligned memory accesses occur when the frst address of a device memory transaction is an even multiple of the cache granularity being used to service the transaction (either 32 bytes for L2 cache or 128 bytes for L1 cache). Performing a misaligned load will cause wasted bandwidth.

Coalesced memory accesses occur when all 32 threads in a warp access a contiguous chunk of memory.

capture

Memory store operations are relatively simple. The L1 cache is not used for store operations on either Fermi or Kepler GPUs, store operations are only cached in the L2 cache before being sent to device memory. Stores are performed at a 32-byte segment granularity. Memory transactions can be one, two, or four segments at a time. For example, if two addresses fall within the same 128-byte region but not within an aligned 64-byte region, one four-segment transaction will be issued (that is, issuing a single four-segment transaction performs better than issuing two one-segment transactions).

capture

CUDA编程笔记(14)——zero-copy memory

这篇笔记摘自Professional CUDA C Programming

In general, the host cannot directly access device variables, and the device cannot directly access host variables. There is one exception to this rule: zero-copy memory. Both the host and device can access zero-copy memory.

GPU threads can directly access zero-copy memory. There are several advantages to using zero-copy memory in CUDA kernels, such as:
➤ Leveraging host memory when there is insufficient device memory
➤ Avoiding explicit data transfer between the host and device
➤ Improving PCIe transfer rates
When using zero-copy memory to share data between the host and device, you must synchronize memory accesses across the host and device. Modifying data in zero-copy memory from both the host and device at the same time will result in undefned behavior.

There are two common categories of heterogeneous computing system architectures: Integrated and discrete.

In integrated architectures, CPUs and GPUs are fused onto a single die and physically share main memory. In this architecture, zero-copy memory is more likely to benefit both performance and programmability because no copies over the PCIe bus are necessary.

For discrete systems with devices connected to the host via PCIe bus, zero-copy memory is advantageous only in special cases.

Because the mapped pinned memory is shared between the host and device, you must synchronize memory accesses to avoid any potential data hazards caused by multiple threads accessing the same memory location without synchronization.

Be careful to not overuse zero-copy memory. Device kernels that read from zero-copy memory can be very slow due to its high-latency.

 

CUDA编程笔记(13)——pinned memory

这篇笔记摘自Professional CUDA C Programming

Allocated host memory is by default pageable, that is, subject to page fault operations that move data in host virtual memory to different physical locations as directed by the operating system. Virtual memory offers the illusion of much more main memory than is physically available, just as the L1 cache offers the illusion of much more on-chip memory than is physically available.

The GPU cannot safely access data in pageable host memory because it has no control over when the host operating system may choose to physically move that data. When transferring data from pageable host memory to device memory, the CUDA driver first allocates temporary page-locked or pinned host memory, copies the source host data to pinned memory, and then transfers the data from pinned memory to device memory, as illustrated on the left side of Figure 4-4.

capture

The CUDA runtime allows you to directly allocate pinned host memory using:
cudaError_t cudaMallocHost(void **devPtr, size_t count);
This function allocates count bytes of host memory that is page-locked and accessible to the device. Since the pinned memory can be accessed directly by the device, it can be read and written with much higher bandwidth than pageable memory. However, allocating excessive amounts of pinned memory might degrade host system performance, since it reduces the amount of pageable memory available to the host system for storing virtual memory data.

 

Unix之殇

前几天网上出现了Solaris项目将会被Oracle停掉的谣言。尽管消息一直未被证实,但是以Solaris为代表的传统Unix操作系统的没落却是不争的事实。在上个月,top500发布的目前世界上运行最快的500台超级计算机中,有498台运行的是Linux

capture

由此可见,LinuxUnix目前的境遇可谓是天壤之别。

我不知道究竟是什么原因造成了目前Linux一统天下的局面,但是可以确定的是一定不是技术领域方面的原因。我没有为IBM公司工作过,也从未接触过AIX操作系统,所以对AIX没有发言权。而对于BSD系列操作系统(FreeBSDOpenBSDNetBSD等等),仅仅限于安装和使用过,并没有什么太深的体会。我为HP/HPE公司效力过,虽然并没有使用过HP-UX,但是周围有很多同事以前是做HP-UX相关工作的:开发新功能,做Unix认证等等。听他们讲,HP-UX非常稳定,很多电信,银行等对稳定性要求特别高的环境仍然在使用着HP-UX,也许这些企业慢慢地会转向Linux?我不知道。。。至于Solaris,我曾经在上面做过4年多的全职开发。Solaris上面有很多很cool的工具供用户使用,比如mdb,比如DTrace,这些工具为我工作提供了巨大的帮助,极大地满足了一个底层软件工程师的好奇心。此外Solaris也是以运行稳定而著称,比如这台已经连续运行了10年的装有Solaris的机器(图片出处:https://pbs.twimg.com/media/CjtxiOmWYAA5lHB.jpg):

cjtxiomwyaa5lhb

再来看看Linux,其实一直以来,Linux系统上并没有可以匹敌DTrace的系统tracing工具,直至最近BPF功能的成熟,可以说在tracing领域落后了Solaris整整12年(可以参考这篇文章:Linux in 2016 catches up to Solaris from 2004);再比如目前Ubuntu发行版中引入的ZFS文件系统,也是出自Solaris。所以,其实如果单单从技术领域来看,Linux不仅不见得做的比Unix好,某些方面甚至还是处于下风的。

在上面提到的几种Unix中,除了BSD系列,其余3种可以说都是某个传统硬件服务器厂商的私有操作系统。虽然曾经有OpenSolaris这个开源产品,但是也仅仅是昙花一现(个人觉得OpenSolaris最大的意义在于由其衍生出了illumos内核,以及基于illumos内核的类Solaris系统。比如smartos。)。所以说,是不是由于最近这些年互联网的日渐强势,硬件厂商的效益江河日下,而“城门失火,殃及池鱼”,随之而来的就是这些Unix也会受到不小的冲击呢?个人觉得应该有一定关系吧。但如果仅仅把Linux成功的原因归结于“开源”,似乎也有失偏颇,BSD系列操作系统也是开源的,且其在license上更为宽松(参考这里:Comparing BSD and Linux)。所以说对于Linux目前具有如此统治力的原因,真的是很难说清。

相信目前很多的中小公司都完全转向Linux了。最直白的原因:人好招。你见过多少招聘信息要求熟悉FreeBSD?肯定没有要求Linux的多。至于要求熟悉NetBSD的?也许有,但是我是没见过。所以对目前Unix人才的需要还是主要在大公司,也只有大公司有意愿和实力做这些“日渐小众”的Unix的相关工作。例如,Brendan Gregg在其社交账号中为对Solaris工程师提到Netflix目前使用FreeBSD

capture2

我很怀念十几年前各大操作系统“百花齐放”的时代,这样想并不是因为我对Linux有任何成见,只是当你的服务器都运行着清一色的Linux操作系统时,实在是觉得有些单调和乏味,就像现在人类使用的手机也基本可以分为两大阵营:iOSAndroid(又是Linux)。世界本应该就是多样化的,丰富多彩的,所以希望其它的Unix有朝一日可以“复兴”吧。。。

nvcc简介

nvcc是“The main wrapper for the NVIDIA CUDA Compiler suite. Used to compile and link both host and gpu code.”,查看其版本可以使用--version选项:

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Sun_Sep__4_22:14:01_CDT_2016
Cuda compilation tools, release 8.0, V8.0.44

关于不同CUDA版本所支持的compute capability可以参考这里

CUDA VERSION   Min CC   Deprecated CC  Default CC
5.5 (and prior) 1.0       N/A             1.0
6.0             1.0       1.0             1.0
6.5             1.1       1.x             2.0
7.0             2.0       N/A             2.0
7.5 (same as 7.0)
8.0             2.0       2.x             2.0

Min CC = minimum compute capability that can be specified to nvcc

Deprecated CC = If you specify this CC, you will get a deprecation message, but compile should still proceed.

Default CC = The architecture that will be targetted if no `-arch` or `-gencode` switches are used

根据CUDA命名规范:GPUs are named sm_xy, where x denotes the GPU generation number, and y the version in that generation.

This situation is different for GPUs, because NVIDIA cannot guarantee binary compatibility without sacrificing regular opportunities for GPU improvements. Rather, as is already conventional in the graphics programming domain, nvcc relies on a two stage compilation model for ensuring application compatibility with future GPU generations.

nvcc可以保证编译出程序的application compatibility,但不能保证binary compatibility。在编译过程中,第一阶段产生virtual GPU architecture code,即PTX;第二阶段才编译出在真实GPU上运行的代码。因此真实的GPU必须实现了virtual GPU所要求的功能。因此,From this it follows that the virtual architecture should always be chosen as low as possible, thereby maximizing the actual GPUs to run on. The real architecture should be chosen as high as possible (assuming that this always generates better code), but this is only possible with knowledge of the actual GPUs on which the application is expected to run.

--gpu-architecture arch编译选项用来指定NVIDIA virtual GPU architecture。例如,compute_50。通常情况下,--gpu-architecture arch是用来生成PTX代码,不会用来生成运行在特定GPU上的代码。--gpu-code code,...选项则是用来指定the name of the NVIDIA GPU to assemble and optimize PTX for。例如,sm_50。关于这两个选项的取值的例子,可以参考这里

参考资料:
NVIDIA CUDA Compiler Driver NVCC
What is the purpose of using multiple “arch” flags in Nvidia’s NVCC compiler?

CUDA编程笔记(12)——CUDA Memory Model

这篇笔记摘自Professional CUDA C Programming

The CUDA memory model exposes many types of programmable memory to you:
➤ Registers
➤ Shared memory
➤ Local memory
➤ Constant memory
➤ Texture memory
➤ Global memory
The following figure illustrates the hierarchy of these memory spaces. Each has a different scope, lifetime, and caching behavior. A thread in a kernel has its own private local memory. A thread block has its own shared memory, visible to all threads in the same thread block, and whose contents persist for the lifetime of the thread block. All threads can access global memory. There are also two read-only memory spaces accessible by all threads: the constant and texture memory spaces. The global, constant, and texture memory spaces are optimized for different uses. Texture memory offers different address modes and filtering for various data layouts. The contents of global, constant, and texture memory have the same lifetime as an application.

capture

capture

The principal traits of the various memory types are summarized in following table:

capture

CUDA编程笔记(11)——warp

这篇笔记摘自Professional CUDA C Programming

Warps are the basic unit of execution in an SM. When you launch a grid of thread blocks, the thread blocks in the grid are distributed among SMs. Once a thread block is scheduled to an SM, threads in the thread block are further partitioned into warps. A warp consists of 32 consecutive threads and all threads in a warp are executed in Single Instruction Multiple Thread (SIMT) fashion; that is, all threads execute the same instruction, and each thread carries out that operation on its own private data. The following figure illustrates the relationship between the logical view and hardware view of a thread block.

capture

From the logical perspective, a thread block is a collection of threads organized in a 1D, 2D, or 3D layout.
From the hardware perspective, a thread block is a 1D collection of warps. Threads in a thread block are organized in a 1D layout, and each set of 32 consecutive threads forms a warp.

在实际的执行中,每个block会被切割成一个一个的warp,而warp中的thread会同步运行。

Threads in the same warp executing different instructions is referred to as warp divergence.
If threads of a warp diverge, the warp serially executes each branch path, disabling threads that do not take that path. Warp divergence can cause signifcantly degraded performance.
Take note that branch divergence occurs only within a warp. Different conditional values in different warps do not cause warp divergence.

warp divergence只会发生在同一个warp中,参考下图:

capture

The local execution context of a warp mainly consists of the following resources:
➤ Program counters
➤ Registers
➤ Shared memory
The execution context of each warp processed by an SM is maintained on-chip during the entire lifetime of the warp. Therefore, switching from one execution context to another has no cost.

Each SM has a set of 32-bit registers stored in a register file that are partitioned among threads, and a fixed amount of shared memory that is partitioned among thread blocks. The number of thread blocks and warps that can simultaneously reside on an SM for a given kernel depends on the number of registers and amount of shared memory available on the SM and required by the kernel.

thread共享registerblock共享shared memory

A thread block is called an active block when compute resources, such as registers and shared memory, have been allocated to it. The warps it contains are called active warps. Active warps can be further classifed into the following three types:
➤ Selected warp
➤ Stalled warp
➤ Eligible warp
The warp schedulers on an SM select active warps on every cycle and dispatch them to execution units. A warp that is actively executing is called a selected warp. If an active warp is ready for execution but not currently executing, it is an eligible warp. If a warp is not ready for execution, it is a stalled warp. A warp is eligible for execution if both of the following two conditions is met:
➤ Thirty-two CUDA cores are available for execution.
➤ All arguments to the current instruction are ready.

GUIDELINES FOR GRID AND BLOCK SIZE
Using these guidelines will help your application scale on current and future devices:
➤ Keep the number of threads per block a multiple of warp size (32).
➤ Avoid small block sizes: Start with at least 128 or 256 threads per block.
➤ Adjust block size up or down according to kernel resource requirements.
➤ Keep the number of blocks much greater than the number of SMs to expose sufficient parallelism to your device.
➤ Conduct experiments to discover the best execution configuration and resource usage.