20200308

@liwei 2020-03-08T00:00:00

Plan

cuda文档阅读
- 3.2.6 多卡系统 40mins
- 3.2.7 统一虚拟地址空间 15mins
- 3.2.8 进程间通信 15mins
leetcode daily

Notes

多卡系统

设备枚举

通常一个host，会配合多个device进行使用，即服务器上插入多张gpu的板卡。我们可以使用固定结构查询一个host上面的所有device，demo程序如下：

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    // capability 为 gpu 对应的架构编号
    printf("Device %d has compute capability %d.%d.\n",
           device, deviceProp.major, deviceProp.minor);
}

设备选择

cuda中的device，是与host上的线程进行绑定的。当host的线程调用了cudaSetDevice()，kernel launch、内存分配、stream等资源将都会与这个device直接关联。一个例子如下：

size_t size = 1024 * sizeof(float);
cudaSetDevice(0);            // Set device 0 as current
float* p0;
cudaMalloc(&p0, size);       // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1);            // Set device 1 as current
float* p1;
cudaMalloc(&p1, size);       // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

当我们在cudaSetDevice()后，一定要注意相应的资源函数，会不会出现失效。

p2p 内存获取与拷贝

在我们的系统中，如果pcie或者nvlink支持，device之间可以互相操作memory。具体系统是否支持这一特性，可以用cudaDeviceCanAccessPeer()这一函数判断。如果支持，用cudaDeviceEnablePeerAccess()打开这一功能。所有的device使用的memory，都在一个统一的虚拟地址空间中。demo程序如下：

cudaSetDevice(0);                   // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size);              // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0);        // Launch kernel on device 0
cudaSetDevice(1);                   // Set device 1 as current
cudaDeviceEnablePeerAccess(0, 0);   // Enable peer-to-peer access
                                    // with device 0

// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);

p2p 内存拷贝

当我们需要在两个device之间进行memcpy等操作时，cuda提供了cudaMemcpyPeer(), cudaMemcpyPeerAsync(), cudaMemcpy3DPeer(), 以及cudaMemcpy3DPeerAsync()接口。使用的时候注意进行同步即可。

统一地址空间

统一地址空间意味着host和所有的device可以通过一个虚拟地址空间进行寻址。使用统一地址空间的必须条件为：

应用地址为64位
device 的 capability 需要在 2.0 以上

这意味着：

任何通过cuda进行分配的地址空间，可以通过cudaPointerGetAttributes()判断其具体的地址
cudaMemcpy*系列函数，cudaMemcpyKind参数可以设置为cudaMemcpyDefault，直接根据指针决定方向
cudaHostAlloc函数分配的空间，可以直接被多卡使用

进程间通信

一个进程中，不同thread创建的device上memory指针或者event handle是互相可见的。如果想要进程之间共享memory或者event信息，我们需要使用api手册中的Inter Process Communication API。只支持64位linux平台。更具体的demo使用时查询手册即可。另外可能需要注意一下的是，不同进程想要通信，需要在相同的driver和runtime上进行编译、链接。

leetcode daily

今天的题目是 322. 零钱兑换

解题过程

初步分析题目，题目应该是动态规划，
递推表达式应该是
1. f(x) = min(f(x - value0) + 1, f(x - value1) + 1, ...)
2. f(x) = 0 when x = 0
3. f(x) = -1 when x < 0
通过画出状态树，可以发现中间状态的f(X)存在较多重复求解，可以进行缓存，以避免超时
因此，在main函数中，单独分配一个全局的vector存储子状态，递归子函数则每次更新和查询这个vector

然而，问题来了，感觉思路毛问题没有，就是死活超时...我的答案：

// f(x) = min(f(x - coins[0]) + 1, f(x - coins[1] + 1, ...))
class Solution {
private:
    vector<int> recorder;
    int dp(vector<int>& coins, int amount) {
        if (amount == 0) return 0;
        if (amount < 0) return -1;
        if (recorder[amount] != -1) return recorder[amount];
        int result = INT_MAX;
        for (auto value : coins) {
            int tmp = dp(coins, amount - value);
            if (tmp >= 0 && tmp + 1 < result) {
                result = tmp + 1;
            }
        }
        if (result == INT_MAX) {
            recorder[amount] = -1;
            return -1;
        } else {
            recorder[amount] = result;
            return result;
        }
    }
public:
    int coinChange(vector<int>& coins, int amount) {
        if (amount == 0) {
            return 0;
        }
        recorder.resize(amount + 1, -1);
        return dp(coins, amount);
    }
};

后来仿照答案，写的可以通过的版本：

// f(x) = min(f(x - coins[0]) + 1, f(x - coins[1] + 1, ...))
class Solution {
    vector<int> count;
    int dp(vector<int>& coins, int rem) {
        if (rem == 0) return 0;
        if (rem < 0) return -1;
        if (count[rem - 1] != 0) return count[rem - 1];
        int MIN = INT_MAX;
        for (auto value : coins) {
            int tmp = dp(coins, rem - value);
            if (tmp >= 0 && tmp < MIN) {
                MIN = tmp + 1;
            }
        }
        count[rem - 1] = MIN == INT_MAX? -1 : MIN;
        return count[rem - 1];
    }
public:
    int coinChange(vector<int>& coins, int amount) {
        if (amount < 1) return 0;
        count.resize(amount);
        return dp(coins, amount);
    }
};

原始答案超时的原因，难道是因为我多写了几个加法么？如果有缘人能看到这篇文章，求个指导...

另一个思路

从头开始建立上面的count数组，直到得到最后的结果。如果说上面的思路是从上到下，这个思路则是从下而上。

总结

动态规划的题目，首先考虑总结状态转移方程以及初始状态函数
树的递归遍历，链表的操作有一个相似的地方是，如果开始状态有一些不同，我们不妨将初始状态独立出来。在链表中加空的头结点，在递归中，原始函数中进行一些初始化，开始调用递归函数。

从今天起，一些非专用的英文单词，统一用小写进行编写。平常工作天天不写代码，时间久了，手真的会非常生疏...practice makes perfect