使用 Grid-Stride Loop 复用 CUDA 线程

在阅读 CUDA 文档的过程中看到这一篇博文:CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops,觉得是很不错的思想(技巧),因此记录下来,并且附上一些验证数据。

In theory

问题起源:一般使用 CUDA 并行计算时,总使用一个线程对应一块数据,例如计算两向量之和:

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
// device code to compute c = a + b;
// this method assumes we have enough threads to do the computation
// enough means larger than the size of array a, b and c
__global__
void addArray(int n, float* a, float* b, float* c)
{
// assume we have 1k elements to compute, in this way
// thread 0 is responsible for c[0] = a[0] + b[0]
// thread 1 is responsible for c[1] = a[1] + b[1]
// and so on
auto index = blockIdx.x * blockDim.x + threadIdx.x;
if (index < n)
c[index] = a[index] + b[index];
}

#define ARRAYSIZE 1000000
#define BLOCKSIZE 1024 // can be any size, better be multiple of 32

__host__
int main()
{
// ...
// calculate the minmal number of blocks to cover all data
auto numBlocks = (ARRAYSIZE + BLOCKSIZE - 1) / BLOCKSIZE;
addArray <<<numBlocks, BLOCKSIZE >>> (ARRAYSIZE, A, B, C);
// ...
}

如我们所知,CUDA 中线程组织结构从高到低分为 Grid、Block、Warp 三层,这三层中,每一层都有自己能容纳的最大线程数量,例如在我的设备(1080Ti)上,每个 Block 最多容纳 1024 个线程,每个 Warp 固定是 32 个线程,而每个 Grid 能容纳的数量就比较大了(2147483647, 65535, 65535)。同时还有一个条件,每个 multiprocessor 最大容纳 2048 个线程。总而言之,GPU 虽然能提供大量的线程,但并不是无限的。

在不讨论显存大小的前提下,总可能出现这样的情况:程序需要的线程数量大于 GPU 可以提供的线程数量。此时上面的代码就不行了,但下面这种处理方式就可以适应:

查看更多

Windows 实现有线 + Wifi 同时上网(分管内外网)

现在的笔记本大多是双网卡,一个有线一个无线。若同时连上有线与 Wifi,Windows 优先使用有线连接。有时我们需要两者皆能联网,例如我的场景:Wifi 连接实验室内网,有线连接校园网,方便挂 PT。那么可按本文操作如下。

首先连接有线与无线,在 CMD 管理员模式中输出当前路由表:

1
route print

顶部两行以 0.0.0.0 开头的路由即分别是有线与无线连接,注意最后一项「跃点数」,数值越低优先级越高。之所以不能在使用有线时同时使用 WiFi 就是由于这两条路由打架,系统按照跃点数选择了有线连接:

1
2
3
网络目标    网络掩码     网关            接口               跃点数
0.0.0.0 0.0.0.0 172.17.104.1 172.17.105.76 26
0.0.0.0 0.0.0.0 192.168.2.1 192.168.2.102 35

查看更多

编译与配置 Boost 库

Boost 是 C++ 的一个增强库,可以看做 std 的扩展。在 Windows 上编译安装很简单。

首先下载代码包:https://www.boost.org/users/history/ 并解压,使用 Visual Studio 提供的 x64 兼容命令提示符工具进入源码目录,运行目录下的 bootstrap.bat 脚本,目录下将出现 bjam.exe 与 b2.exe 等文件。之后即可使用 bjam 工具编译。

实际上 Boost 库的绝大多数组件都是 header-only 的,也就是无需进一步编译生成可执行文件,只需要引入头文件即可,但是以下几个组件除外:

  • Boost.Chrono
查看更多

使用 PHP 执行 shell 脚本的权限问题

上一篇讲使用 GitHub webhook 来部署的问题,其中一步是用 PHP 执行一段 shell 脚本。类似于如下的 PHP 代码:

1
2
<?php
shell_exec('cd somepath && git pull --rebase');

把这段代码保存为 test.php,在终端中使用 php test.php 执行可以得到正确的结果,但是在浏览器中访问这个 PHP 却不行。

除了安全模式、shell_exec 被禁用的问题之外,最大的可能是运行权限的问题。尝试在终端与浏览器中访问包含以下代码的 PHP 文件:

1
2
<?php
shell_exec('whoami');

查看更多

使用 GitHub webhook 自动部署

我的三个静态站点:三無計劃無知識無項目均使用 Travis-CI 自动构建。在此之前,三个站点都托管在 GitHub Pages 或者 Coding Pages 上,但前者在大陆速度缓慢,后者时常抽风,总体表现都不如我自己的服务器,因此决定将站点放在自己这里。

GitHub webhook 可在仓库发生某些事件(例如 push)时向某指定 URL 发起 POST 请求,请求 body 就是详细的 push 事件内容,因此可以在服务器上根据请求部署站点,比如直接从 Github 拉取最新源码然后构建,或者由 CI 服务构建好之后拉取最新的 build 到服务器上。


首先保存以下 PHP 到服务器上可从外部访问的地方,并在 Github 对应仓库设置中添加一个 hook,URL 就填该 PHP 的网址,例如 https://api.imalan.cn/somepath/hook.php

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
<?php
/**
* GitHub Webhook 服务端处理
* 根据 Repo 与 Branch 的 push 事件运行脚本
*
* @author 熊猫小A | AlanDecode
* @link https://www.imalan.cn
*/

function run($data, $config) {
foreach ($config['endpoints'] as $endpoint) {
// check repository and branch
if ($data['repository']['full_name'] != $endpoint['repo']) {
continue;
}
if ($data['ref'] != 'refs/heads/' . $endpoint['branch']) {
continue;
}

// excute
try {
echo shell_exec($endpoint['run']);
echo $endpoint['action'];
} catch (Exception $e) {
echo $e->getMessage();
}
}
}

try {
// data
$rawInput = file_get_contents('php://input');
$data = json_decode($rawInput, true);

// config
$config_path = __DIR__.'/config.json';
if (!file_exists($config_path)) {
throw new Exception("Can't find ".$config_path);
}
$config = json_decode(file_get_contents($config_path), true);

// verify
if (!empty($config['key'])) {
if (!isset($_SERVER['HTTP_X_HUB_SIGNATURE'])) {
die('No signature presented.');
}

list($algo, $hash) = explode('=', $_SERVER['HTTP_X_HUB_SIGNATURE'], 2);
$signature = hash_hmac($algo, $rawInput, $config['key']);

if ($signature != $hash) {
die("Signature doesn't match.");
}
}

// excute
run($data, $config);

} catch (Exception $e) {
echo $e->getMessage();
}

查看更多

Git 代理配置方案

国内连 Github 是个很头疼的问题。浏览器还好,只要电脑挂了代理一般就没什么问题;主要是终端使用有些麻烦。折腾出来的最佳实践如下。

部分来自:https://molunerfinn.com/git-ssh2https/

设置代理

首先,代码仓库尽量使用 HTTPS 协议,也就是以 https://github.com 开头的链接,而不是以 `git@github.com开头的 SSH 链接,这样就可以通过git config` 方便地设置代理:

1
git config --global http.proxy 'socks5://127.0.0.1:1080'

网上广泛流传的针对 https.proxy 也设置一遍代理是不必要的。

查看更多

常用设置代理命令

本文整理了 Windows 命令行 和 Linux 终端中设置代理的命令。以本地 HTTP/HTTPS 代理 127.0.0.1:8888 和 SOCKS5 代理 127.0.0.1:8889 为例。

Windows 命令行代理设置

HTTP 代理设置:

1
2
set http_proxy=http://127.0.0.1:8888
set https_proxy=http://127.0.0.1:8888

SOCKS5 代理设置:

1
2
set http_proxy=socks5://127.0.0.1:8889
set https_proxy=socks5://127.0.0.1:8889

查看更多

C++智能指针:原理与实现

若问起 Java 与 C++ 在使用体验上的差别,许多人都会提到「垃圾回收」这个词,Java 自带垃圾回收机制,而 C++ 没有。同样是面向对象的程序语言,这一点上的设计理念却如此大相径庭。有个古老的段子说,Java 的设计者认为:“内存管理这么重要的事情怎么能交给愚蠢的程序员呢!”,而 C++ 设计者认为:“内存管理这么重要的事情怎么能交给愚蠢的机器呢!”二者好像都有些道理。

C/C++ 的指针把计算机底层暴露给了程序员,在使其变得灵活强大的同时也增加了程序员的负担,特别是对初学者来说,内存忘记释放、野指针满天飞都是常见问题,而且没那么容易克服。即使是对有经验的程序员来说,随着项目规模变大,资源管理也会成为一个令人头疼的问题。

为了解决内存泄漏的难题,C++ 11 标准引入了三种智能指针:std::shared_ptrstd::unique_ptrstd::weak_ptr。但只会用是不足够的,这篇文章就来说说智能指针的基本思想与实现方法。注意,文章中的代码自然与 std 的实现有所不同,旨在抓住思想本质,免去为了适应特殊情况而引入的更多细节。若实际项目中真遇到那些所谓「特殊情况」,则仔细研读 std 的代码是很好的方案。

裸指针有什么问题?

首先要明确的是我们究竟想要解决什么问题。在 C++ 中通过 new 关键字在堆上申请空间后,除非使用 delete 关键字释放这块资源,否则这块资源对系统来说就是被占用的,直到整个程序结束才会被回收。即使当前程序已经不再需要这块资源,或者已经丧失了对这块资源的控制,也是如此。一个简单的例子:

1
2
for(int i=0; i<100; i++)
char* p = new char[100];

查看更多