我的站点

一个系统软件工程师的随手涂鸦

Page 2 of 70

CUDA编程笔记(17)——Matrix transpose (shared memory)

An Efficient Matrix Transpose in CUDA C/C++Coalesced Transpose Via Shared Memory一节讲述如何使用shared memory高效地实现matrix transpose

__global__ void transposeCoalesced(float *odata, const float *idata)
{
  __shared__ float tile[TILE_DIM][TILE_DIM];

  int x = blockIdx.x * TILE_DIM + threadIdx.x;
  int y = blockIdx.y * TILE_DIM + threadIdx.y;
  int width = gridDim.x * TILE_DIM;

  for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
 tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];

  __syncthreads();

  x = blockIdx.y * TILE_DIM + threadIdx.x;  // transpose block offset
  y = blockIdx.x * TILE_DIM + threadIdx.y;

  for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
 odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
}

(1)idataodata分别是表示1024X1024float元素的matrix的连续内存:

IMG_20170216_145959[1]

(2)关于blockIdxthreadIdx的取值,参考下面的图:

IMG_20170216_151045[1]

shared memory请参考下面的图:

IMG_20170216_151058[1]

(3)在下列代码中:

for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
    tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];

每一个block32X8大小,需要循环4次,把一个block内容copytile这个shared memory中。idata是按行读取的,因此是coalesced

IMG_20170216_152636[1]

(4)最难理解的在最后一部分:

x = blockIdx.y * TILE_DIM + threadIdx.x;  // transpose block offset
y = blockIdx.x * TILE_DIM + threadIdx.y;

for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
    odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];

对比从idata读取数据和写数据到odata

......
tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
......
odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
......

可以看到是把tile做了transpose的数据(行变列,列变行)传给odata。而确定需要把tile放到哪里位置的代码:

x = blockIdx.y * TILE_DIM + threadIdx.x;  // transpose block offset
y = blockIdx.x * TILE_DIM + threadIdx.y;

假设blockIdx.x31blockIdx.y0threadIdx.x1threadIdx.y2。根据上述代码,计算xy

x = 0 * 32 + 1;
y = 31 * 32 + 2;

根据下面的图,可以看到是把东北角的内容copy的西南角:

IMG_20170216_155616[1]

与*NIX有关的杂志

本文介绍一些我接触过的与*nix有关的杂志。

首先要提到的就是Linux Journal(官方网址:http://www.linuxjournal.com/)。根据Wikipedia的介绍,这应该是最早的一本介绍Linux的杂志:

Linux Journal was the first magazine to be published about the Linux kernel and operating systems based on it. It was established in 1994.

不过Linux Journal现在不再发行纸质版了,只提供电子版。该杂志这段时间有一项促销活动:即截至到今年328日前,你只需花28.5美元(使用优惠码:2017ARCH可免10美元),就可以购买到从1994年到2016年杂志的电子合订版:

1

个人觉得还是很划算的,虽然有些文章已经年代久远,但是还是很有参考价值的。另外,Linux Journal也会将其文章发布到官网上供读者免费阅读。因次,是否愿意花钱买合订本或者订阅,就“仁者见仁,智者见智”了。

再说一下Linux Format(官方网址:http://www.linuxformat.com/)和Linux Voice(官方网址:https://www.linuxvoice.com/),二者都是英国出版的Linux杂志。不知是否因为Linux Voice的主创团队均出自Linux Format的缘故,二者有太多类似之处:都同时提供纸质版和电子版,都提供过刊的免费下载,等等。也许是由于版权原因,我在国内没见到过这两种杂志。在国外我阅读过纸质版,制作很精美,每期还附赠光盘。感兴趣的朋友可以下载它们提供的过刊了解一下。

以上提到的都是以Linux为主题的杂志,再介绍一个以BSD为主打内容的杂志:BSD magazine(官方网址:https://bsdmag.org/)。这是一本真正免费的杂志,订阅以后,你不需花一分钱,就会收到每一期。从这本杂志里,你可以获悉时下BSD家族的最新动态,虽然偶尔也会有Linux的内容出现。去年年底,这本杂志一度宣布要停刊了,不过目前又撤销了这个决定。我个人很希望这本杂志可以继续办下去。

最后,我很希望自己的国家能有一本中文版的Linux杂志。但是想想现在的情形,我们的时间都被其它的事物占去了,也许根本无法诞生这样一本杂志了。。。

后记:对于其它的类似杂志,比如:http://www.linux-magazine.com/。因为我没有一点了解,就不发表评论了。

git小技巧(10)——如何同original master进行rebase

original repositorymaster branch rebase的方法:
(1)

git fetch origin # Updates origin/master
git rebase origin/master # Rebases current branch onto origin/master

(2)

git pull --rebase origin master

参考资料:
How to rebase local branch with remote master

2017年1月总结

工作方面:
a)这个月是在调试中度过的:两周调试多线程bug(参考这篇文章:An example of debugging parallel program),一周调试性能问题。收获颇丰:-);
b)学习了OpenMP技术。

业余时间(技术):

a)学习了这套Linux X86_64汇编教程
b)写了3Unix公众号文章。

业余时间(读书):
Clean, Well-Lighted Sentences: A Guide to Avoiding the Most Common Errors in Grammar and Punctuation :这本书真得很不错,很薄。利用小长假时间读了一下,弄清楚很多英语语法疑问。当然一遍不一定都能记住,还需要反复阅读。

如何选择Linux发行版?

@YeimMeMeMes 在其个人社交媒体主页上贴出了一张如何选择Linux发行版的图(原图链接:https://pbs.twimg.com/media/C3PoAagWMAAe5YT.jpg:large):

C3PoAagWMAAe5YT

因为这幅图上所列举的大部分Linux发行版我都没怎么使用过,所以我没有资格对这张图是否合理发表评论,仅供读者参考。下面我仅对我曾经使用过的一些Linux发行版谈些感受:

(1)这个图上并没有出现一些Linux的商业版本,譬如,RedHatRed Hat Enterprise LinuxSuSESUSE Linux Enterprise Server,等等。我以前做过这两个Linux商业版本上的测试工作,感觉还是很稳定的。给我一个很深的印象就是在Red Hat Enterprise Linux上编译和安装最新版本的Linux内核总是非常顺利,而在其它发行版上有时会遇到一些莫名其妙的问题。从去年开始,Red Hat Enterprise Linux对开发者已经免费了(可参考这篇文章:As in beer: Red Hat offers RHEL free to developers),有兴趣的朋友可以体验一下。

(2)最近两个月,在工作中我主要使用Arch LinuxArch Linux在软件包更新方面非常及时。有时你发现当前软件包不是最新的,只要在其网站上提一个请求,很快就会有维护者响应。如果你总是希望可以使用上最新的软件版本,不妨尝试一下Arch Linux

Linux系统上如何查看进程(线程)所运行的CPU

本文介绍如何在Linux系统上查看某个进程(线程)所运行的CPU,但在此之前我们需要弄清楚两个基本概念:

(1)Linux操作系统上的进程和线程没有本质区别,在内核看来都是一个task。属于同一个进程的各个线程共享某些资源,每一个线程都有一个ID,而“主线程”的线程ID同进程ID,也就是我们常说的PID是一样的。

(2)使用lscpu命令,可以得到当前系统CPU的数量:

$ lscpu
......
CPU(s):                24
On-line CPU(s) list:   0-23
Thread(s) per core:    2
Core(s) per socket:    6
Socket(s):             2
......

系统有2个物理CPUSocket(s): 2),每个CPU6coreCore(s) per socket: 6),而每个core又有2hardware threadThread(s) per core: 2)。所以整个系统上一共有2X6X2=24CPU(s):24)个逻辑CPU,也就是实际运行程序的CPU

使用htop命令可以得到进程(线程)所运行的CPU信息,但是htop默认情况下不会显示这一信息:

1
开启方法如下:
(1)启动htop后,按F2Setup):

2
(2) Setup中选择Columns,然后在Available Columns中选择PROCESSOR - ID of the CPU the process last executed, 接下来按F5Add)和F10Done)即可:

3

现在htop就会显示CPU的相关信息了。需要注意的是,其实htop显示的只是“进程(线程)之前所运行的CPU”,而不是“进程(线程)当前所运行的CPU”,因为有可能在htop显示的同时,操作系统已经把进程(线程)调度到其它CPU上运行了。

下面是一个运行时会包含4个线程的程序:

#include <omp.h>

int main(void){

        #pragma omp parallel num_threads(4)
        for(;;)
        {
        }

        return 0;
}

编译并运行代码:

$ gcc -fopenmp thread.c
$ ./a.out &
[1] 17235

使用htop命令可以得到各个线程ID,以及在哪个CPU上运行:

4

参考资料:
How to find out which CPU core a process is running on
闲侃CPU(一)

Linux线程模型浅析

Linux的线程是“轻量级进程”(Light-Weight Process,即LWP)。在Linux系统上运行一个程序时,操作系统会为这个程序创建一个进程,其实也就是“主线程”,后续则可以产生出更多的线程。每个进程都有一个PIDProcess ID),每个线程也会有一个TIDThread ID),属于同一进程的线程各自有拥有不同的TID,但它们的PID是相同的,都等于“主线程”的TID。因此从本质上来讲,Linux系统下的进程和线程没有区别,只不过同一进程中的线程可以共享某些资源。下面看一个例子:

#include <unistd.h>
#include <omp.h>

int main(void){

        #pragma omp parallel num_threads(4)
        for(;;)
        {
            sleep(1);
        }

        return 0;
}

编译并在后台运行这个程序:

$ gcc -fopenmp threads.c
$ ./a.out &
[1] 9802

进程的PID9802,用ps -T pid命令查看进程的线程信息:

$ ps -T 9802
  PID  SPID TTY      STAT   TIME COMMAND
 9802  9802 pts/1    Sl     0:00 ./a.out
 9802  9803 pts/1    Sl     0:00 ./a.out
 9802  9804 pts/1    Sl     0:00 ./a.out
 9802  9805 pts/1    Sl     0:00 ./a.out

其中SPID即为TID。可以看到当前进程的PID9802,共包含4个线程,其TID依次为:9802980398049805,其中PIDSPID相同的线程即为主线程。

CUDA编程笔记(16)——Shared Memory

这篇笔记摘自Professional CUDA C Programming

Global memory is large, on-board memory and is characterized by relatively high latencies. Shared memory is smaller, low-latency on-chip memory that offers much higher bandwidth than global memory. You can think of it as a program-managed cache. Shared memory is generally useful as:
➤ An intra-block thread communication channel
➤ A program-managed cache for global memory data
➤ Scratch pad memory for transforming data to improve global memory access patterns

Shared memory is partitioned among all resident thread blocks on an SM; therefore, shared memory is a critical resource that limits device parallelism. The more shared memory used by a kernel, the fewer possible concurrently active thread blocks.

 

2016年终总结

2016年就这样过去了,它真的是令我难忘的一年。在这一年中我经历了太多的“第一次”,也有了很多新收获。在2017年的第一天,我再好好地回忆一下这“跌宕起伏”的一年吧。

首先,我在自己的个人博客上开辟了“每月简讯”这个专栏,用来总结每个月的经历。这样可以更清楚地了解自己过去一个月的状态,知道该从哪些方面改进和提高自己。

1月份的时候,我开通了一个微信新公众号:Unix。开始的时候,是每天发一个小tip,后来感觉这种形式其实对人们的帮助并不大。所以在停办一段时间后,又走到发原创文章这条路上。目前发表了不到10篇作品,不敢说写的多好,只能说都是自己用心之作。

DTrace这个公众号这一年并没有写什么的有价值的文章,主要是现在工作中基本不会用到DTrace,所以也没有什么好的材料用来分享。

7月份参加了人生第一次的IELTS考试。

今年写了两个初级教程:一个是关于Go语言:Go 101 Hacks;一个是关于FreeBSDFreeBSD 101 Hacks

1月份至9月份,我在H公司工作。工作主要侧重在两部分:Docker性能测试和Swarm/Swarmkit的功能开发。很遗憾,9月份的时候,公司结构调整,我们整个部门被裁掉了。但是H公司还是很厚道的,给了足够的补偿。在这个时候,才能体现一家公司是否真的是“人性化”。

9月中旬到12月份,我一直处于失业状态,也是我第一次失业。整天过得浑浑噩噩。以前每天上班很辛苦,但是并不觉得有多累。现在闲下来了,反而每天无精打采,还生了病。看来人真的是要折腾的。

12月份来到了当前的公司,开始了新的工作。目前感觉还好,没有什么不适应。

中英文博客这一年中还是坚持更新。中文博客更多的是记一些笔记,英文博客倒是写了一些个人还比较满意的文章。

年初的时候第一次出国,去了新加坡,感受了一下“异国他乡”是什么样子。

好了,就这样吧。新年新开始!

2016年12月简讯

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

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

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

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

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

Page 2 of 70

Powered by WordPress & Theme by Anders Norén