一点说明

cuda与mgpu 概念的对应

以下讨论一维grid的情况下cuda和mgpu的对应关系。

cuda 从 gridDim 和 blockDim 的维度描述一个grid。 mgpu关注线程，一个block的线程数，即blockDim.x，在mgpu中表示为nt，block个数，即gridDim，表示为num_ctas. mgpu启动kernel时给定线程总数count和 nt。因此 num_ctas = div_up(count, nt).

iterate

stride_iterate(func_t, int tid, int count)

count: grid的线程总数

tid: threadIdx.x

blockDim.x 配置的常数，来自lauch_box

gridDim.x = (count + blockDim.x - 1)/blockDim.x

等价于：

 for(int i = 0; i < vt; ++i)
 {
    j = nt * i + tid;
    f(i,j);
 } 

如果忽略参数i，即stride的第几轮，那么等价于

int tid = threadIdx.x;
int stride = blockDim.x;
for(int i=tid; i<count; i+=stride)
{
    f(i);
}

MGPU假设grid宽度足够大，有足够的block覆盖输入数据。即一个block只处理一片数据。所以stride 不是常用的 blockDim.x * gridDim.x, tid 也不是常用的 blockDim.x * blockIdx.x + threadIdx.x

transform(lambda, count, context)

启动count个线程，每个线程执行一次lambda 对应代码如下：

__global__ void kernel(args...)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if(tid < count) lambda( tid, args... );
}

启动方式：

block = nt;
grid = (count + nt - 1) / nt;
kernel<<<grid, block>>>(args...);