1. 介绍

1.1. 什么是CUB

CUB是一个只包含头文件的C++库，可为MXMACA（MetaX Advanced Compute Architecture）编程模型的每一层提供最先进的、可重用的软件组件，包括并行原语和许多实用程序。并行原语包括线程束（warp）级别的“集合”原语、块级别的“集合”原语和设备级别的原语。这些实用程序包括复杂的迭代器、线程和线程块 I/O、设备、内核和存储管理。

1.2. 安装

安装MXMACA工具包将把CUB和Thrust的头文件复制到系统的标准MXMACA include 目录中。

# header location
${MACA_PATH}/include/cub
${MACA_PATH}/include/thrust

不需要单独构建CUB。要在代码中使用CUB原语，只需：

获取最新的CUB发行版
在 MXMACA C++ 源代码中，通过#include指令包含 cub/cub.cuh 这个“伞”（master）头文件。（或者，通过#include指令包含特定的头文件，用来定义需要的 CUB 原语。）
使用MXMACA的mxcc编译器编译您的程序，指定一个 -I<path-to-CUB> include-path 标志，以引用CUB和Thrust头文件库的位置。

1.3. 在应用程序中使用CUB

以下的代码段说明了一个MXMACA内核，其中每块 BLOCK_THREADS 个线程将共同加载、排序和存储自己（BLOCK_THREADS * ITEMS_PER_THREAD）个整数键：

// myCubApp1.cpp
#include <mc_runtime.h>
#include <cub/cub.cuh>
//
// 块排序MXMACA核函数
//
template <int BLOCK_THREADS, int ITEMS_PER_THREAD>
__global__ void BlockSortKernel(int *d_in, int *d_out)
{
    // 专门定义BlockLoad、BlockStore和BlockRadixSort集合类型
    typedef cub::BlockLoad<
        int, BLOCK_THREADS, ITEMS_PER_THREAD, cub::BLOCK_LOAD_TRANSPOSE> BlockLoadT;
    typedef cub::BlockStore<
        int, BLOCK_THREADS, ITEMS_PER_THREAD, cub::BLOCK_STORE_TRANSPOSE> BlockStoreT;
    typedef cub::BlockRadixSort<
        int, BLOCK_THREADS, ITEMS_PER_THREAD> BlockRadixSortT;
    // 为集合操作分配类型安全、可重复使用的共享内存
    __shared__ union {
        typename BlockLoadT::TempStorage       load;
        typename BlockStoreT::TempStorage      store;
        typename BlockRadixSortT::TempStorage  sort;
    } temp_storage;
    // 获取此块在连续键中的段（按线程分块）
    int thread_keys[ITEMS_PER_THREAD];
    int block_offset = blockIdx.x * (BLOCK_THREADS * ITEMS_PER_THREAD);
    BlockLoadT(temp_storage.load).Load(d_in + block_offset, thread_keys);

    __syncthreads();    // 共享内存重用的障碍
    // 多个线程协同对键进行排序
    BlockRadixSortT(temp_storage.sort).Sort(thread_keys);
    __syncthreads();    // 共享内存重用的障碍
    // 存储已排序的段
    BlockStoreT(temp_storage.store).Store(d_out + block_offset, thread_keys);
}

int main()
{
    // 在主程序的其他位置：对块排序操作进行参数化设置并启动该操作
    // 内核中的每个块由128个线程组成，每个块对256个键的段进行排序
    constexpr int TOTAL_ITEMS = 256;
    constexpr int NUM_BLOCKS = 1;
    constexpr int BLOCK_THREADS = 128;
    constexpr int ITEMS_PER_THREAD = TOTAL_ITEMS / NUM_BLOCKS / BLOCK_THREADS;
    int h_in[TOTAL_ITEMS];
    int h_out[TOTAL_ITEMS];
    printf("The items in array h_in:\n");
    for (int i = 0; i < TOTAL_ITEMS; ++i) {
        h_in[i] = TOTAL_ITEMS - i; // [256, 255, 254, ..., 1]
        printf("h_in[%d]: %d\n", i, h_in[i]);
    }
    int *d_in;
    int *d_out;
    mcMalloc((void**)&d_in, sizeof(int) * TOTAL_ITEMS);
    mcMemcpy(d_in, h_in, sizeof(int) * TOTAL_ITEMS, mcMemcpyHostToDevice);
    mcMalloc((void**)&d_out, sizeof(int) * TOTAL_ITEMS);
    BlockSortKernel<BLOCK_THREADS, ITEMS_PER_THREAD><<<NUM_BLOCKS, BLOCK_THREADS>>>(d_in, d_out);
    mcMemcpy(h_out, d_out, sizeof(int) * TOTAL_ITEMS, mcMemcpyDeviceToHost);
    printf("The items in array h_out:\n");
    for (int i = 0; i < TOTAL_ITEMS; ++i) {
        printf("h_out[%d]: %d\n", i, h_out[i]); // [1, 2, 3, ..., 256]
    }
    mcFree(d_in);
    mcFree(d_out);
    return 0;
}

在linux上编译上述应用程序，可以使用以下命令：

export MACA_PATH=your/maca/toolkits/path
export MACA_CLANG_PATH=${MACA_PATH}/mxgpu_llvm/bin
export LD_LIBRARY_PATH=${MACA_PATH}/lib:$LD_LIBRARY_PATH

${MACA_CLANG_PATH}/mxcc myCubApp1.cpp -o myCubApp1 -I ${MACA_PATH}/include/cub

在这个例子中，线程使用cub::BlockLoad、cub::BlockRadixSort和cub::BlockStore来共同加载、排序和存储块的输入项段。因为这些操作是并行的，所以每个原语都需要分配共享内存，以便线程进行通信。

一个CUB集合的典型使用模式是：针对手头的特定问题设置静态专门化原语，例如，正在排序的数据类型、每个块的线程数、每个线程的键数，可选的算法替代方案等。（CUB原语也被目标编译架构隐式特化。）在共享内存空间中分配（或别名）一个特化原语的嵌套TempStorage类型的实例。指定通信细节（例如，分配TempStorage内存）来构造该原语的实例。调用原语实例上的方法。

具体来说，cub::BlockRadixSort用于对分配到线程块的数据项段进行集体排序。为了提供对设备内存的合并访问，我们配置了cub::BlockLoad和cub::BlockStore原语，以便使用条带化访问模式（连续的线程同时访问连续的项）访问内存，然后将键转换为跨线程的块状排列。为了在所有三种原语之间重复使用共享内存，线程块会静态分配它们的TempStorage类型的联合体。