Install

1
2
3
4

# cargo
curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh
# zellij
cargo install --locked zellij

Commands

Open Session

1

zellij -s {sessioin-name} # open new session with name

Attach

1

zellij a {session-name}

Tips

Copy Contest

1
2
3
4

# enter search mode -> edit
Ctrl+s -> e
# change editor
scrollback_editor "nvim"

Troubleshotting

Shortkey Conflict

1
2

# enter locked mode
Ctrl+g

Troubleshotting

unknown terminal type

报错信息

1

'xterm-kitty': unknown terminal type.

解决方案

1
2
3

# append this into shell config file
alias ssh="kitty +kitten ssh" # 但是这样在其他终端没法使用 ssh 甚至某些 ide 无法克隆东西
alias kssh="kitty +kitten ssh"

AQL

官方文档。

Insert

1

INSERT INTO test.demo (PK, username, email) VALUES ('user1', 'John', 'john@example.com');

Select

1

SELECT * FROM test.demo WHERE PK = 'user1';

Arrow

Arrow C++ 库由多个部分组成，每个部分完成不同的功能。

• 物理层
• 一维层
• 二维层
• 计算层
• IO 层
• IPC 层（Inter-Process Communication）
• 格式层
• 设备层
• 文件系统层

物理层（The physical layer）

• Memory management 抽象提供了统一的 API 来管理内存

• Buffer 抽象表示物理数据的连续区域。

一维层（The one-dimensional layer）

• Data Types 赋予物理数据逻辑含义
• Arrays 装了一个或多个有类型的缓冲区，允许将他们视为逻辑连续的值序列（可能是嵌套的）
• Chunked arrays （分块数组）是对数组的泛化，将多个相同类型的数组组合成一个更长的逻辑值序列

二维层（The two-dimensional layer）

• Schemas 描述了多个数据的逻辑集合，每个数据都有不同的名称和类型，以及可选的元数据
• Tables 是根据 Schemas 的分块数组的集合。它们是 Arrow 中最有能力提供数据集的抽象
• Record batches 是由 Schemas 描述的连续数组的集合，允许表的增量构造或序列化

计算层（The compute layer）

• Datums 是灵活的数据集一样，比如可以保存一个 Array 或 Table 引用
• Kernels 是在一组给定的 Datums 上循环运行的专用计算函数，这些数据表示函数的输入和输出参数
• Acero 是一个流执行引擎，允许将计算表示为可以转换数据流的运算符图

IO 层（The IO layer）

• Streams 允许对各种外部数据（例如压缩或内存映射）进行无类型的顺序或可查找访问

IPC 层（The Inter-Process Communication (IPC) layer）

• Messaging Format 允许 Arrow 数据在进程间交换，并尽可能少得拷贝

文件格式层（The file formats layer）

读写尽可能多的文件格式，比如 Parquet、CSV、Orc，以及 Arrow 专用的 Feather 格式。

设备层（The devices layer）

• CUDA 提供基础集成，允许描述由 GPU 分配的内存数据

文件系统层（The filesystem layer）

文件系统抽象允许从不同的存储后端读取和写入数据，例如本地文件系统或 S3 存储桶

计算

通用计算 API

Functions 及其注册

• 函数表示对可能不同类型的输入的计算操作
• 在内部，一个函数由一个或多个“内核”（Kernels）实现
• 函数存储在全局函数注册表（FunctionRegistry ）中，可以在其中按名称查找它们

Input Shapes

• 计算 API 的输入被标识为通用的 Datum 类，可以标识 Scalar、Array、ChunkedArray
• 很多函数支持 Array、Chunked Array和 Scalar 作为输入，但是有些需要特定的输入
  • array_sort_indices，仅接受一个 Array 作为输入
  • sort_indices，更通用，接受 Array、Chunked Array、BatchRecord、Table

调用函数

• 计算函数可以使用 arrow::compute::CallFunctioni() 来调用
• 也可以用具体的 API，比如 arrow::compute::Add()

1

arrow::compute::CallFunctioni("add", ...) 等价于 arrow::compute::Add(...)

需要注意的是，一些函数还需要特定类型的参数（Options）。

1
2

ScalarAggregateOptions scalar_aggregate_options;
arrow::compute::CallFunction("min_max", {...}, &scalar_aggregate_options)

隐式转换（Implicit casts）

functions 可能会对输入进行隐式转换。

可用的 Functions

类型分类

• Numeric，Integer、Floating、Decimal128/256
• Temporal，Date、Time、Timestamp、Duration、Interval
• Binary Like，Binary、LargeBinary、FixedSizeBinary
• String Like，String、LargeString
• List Like，List、LargeList、ListView、LargeListView、FixedSizeList
• Nested，List-likes、Struct、Union、related types like Map

如果一个 Function 不支持一个输入类型，会返回 TypeError Status。

Aggregations（聚合）

标量聚合操作操作在 （Chunked）Array、Scalar，产生单个输出值。

具体的聚合方法看这里。

Group Aggregations（组聚合）

组聚合方法通常无法直接调用，通常作为 SQL 类型 group by 操作，详见这里。

Element-wise (“scalar”) functions

这类方法接受 Array 和 Scalar 作为输入。对于一元输入，输入是什么类型，输出就是什么类型。

两元输入时有如下的语义（在其他领域有时称为 broadcasting 比如 Numpy）。

• (Scalar, Scalar)，输出 Scalar
• (Array, Array)，输出 Array
• (Scalar, Array)、(Array, Scalar) 输出 Array。标量输入被处理为与另一个相同长度的数组，元素是重复的值。

更多。

Acero

Acero 是一个流查询引擎，用来制定和执行计算。

Arrow 使用 ExecPlan 来表示计算，ExecPlan 以零个或多个 Stream 作为输入，但输出只有一个 Stream。Plan 描述了数据怎么被转换、传递。

Acero 处理流式数据，Compute 处理内存中的数据。

概念

ExecNode

ExecNode 是 Acero 中最基础的概念，ExecNode 有零个或多个输入，零或一个输出。如果没有 Source，称其为源；如果没有输出，称其为 Sink。ExecNode 还有多种类型。

• Sources Nodes
• Compute Nodes
• Arrangement NOdes
• Sink Nodes

ExecBatch

批量数据（Batches of Data）使用 ExecBatch 类来描述。ExecBatch 是一个二维结构，和 RecordBatch 很相似。可以有零个或多个列，每个列必须有相同的长度。

1
2

exec_batch.length :        数据行度 (row)
exec_batch.values.size() : 列数量 (column)

RecordBatch 和 ExecBatch 对 Arrays 和 Buffers 都有强所有权。把 RecordBatch 转换为 ExecBatch 总是零拷贝的。但是 ExecBatch 转换为 RecordBatch 时只有在无标量（Scalar） 时才是零拷贝。

ExecPlan

ExecPlan 表示由 ExecNode 组成的图。一个合法的 ExecPlan 必须由一个 Source 节点，但技术上不需要有 Sink 节点。ExecPlan 包含所有节点共享的资源，以及工具方法来启动、停止节点的执行。ExecPlan 和 ExecNode 都与单个执行的生命周期相关联。它们有状态，预计不可重新启动。

Declaration

Declaration 是 ExecNode 的蓝图。声明可以组合成一个图表，形成 ExecPlan 的蓝图。

使用

Acero 的基础工作流如下。

1. 首先，创建 Declaration 图，描述 Plan
2. 调用一个 DeclarationToXyz 方法来执行 Declaration
   1. 创建 ExecPlan。一个新的 ExecPlan 会从 Declaration Graph 中被创建。每个 Declaration 对应 Plan 中的一个 ExecNode。此外，会根据具体的 DeclarationToXyz 创建 Sink Node
   2. 执行 ExecPlan。通常，这是 DeclarationToXyz 调用的一部分，但在 DeclarationToReader 中，阅读器会在计划完成执行之前返回。
   3. 销毁。Plan 一但结束就会被销毁。

创建 Plan

使用 Substrait

Substrait 是创建 Plan （Graph of Declaration）的推荐机制，原因如下。

• Substrait 耗费了大量时间和精力创建用户友好的 API，以简单的方式创建负责的 Execution Plans
• 更容易迁移到其他类似引擎
• 更多基于 Substrait 的工具

程序化构建 Plan

以编程方式创建执行计划比从 Substra 创建计划更简单。Declaration::Sequence() 可以更方便地创建处理序列。

执行 plan

• DeclarationToTable，累计所有的结果到一个 arrow::Table
• DeclarationToReader，允许以迭代的方式使用所有结果，会创建一个 arrow::RecordBatchReader
• DeclarationToStatus，如果只想执行 Plan 而不消费结果时很有用
• 直接运行 Plan
  • 创建 ExecPlan
  • 添加 Sink Node 到 Declaration Graph 中
  • 使用 Declaration::AddToPlan 添加 Declaration 到 Plan 中
  • 验证 Plan，ExecPlan::Validate
  • 开始 Plan，EexcPlan::StartProducing
  • 等待 Future，ExecPlan::finished

可用的 ExecNode

Sources Node

用作数据源。

• source，SourceNodeOptions，通用源节点，包装异步流数据
• table_source，TableSourceNodeOptions
• record_batch_source，RecordBatchSourceNodeOptions
• record_batch_reader_source，RecordBatchReaderSourceNodeOptions
• exec_batch_source，ExecBatchSourceNodeOptions
• array_vector_source，ArrayVectorSourceNodeOptions
• scan，arrow::dataset::ScanNodeOptions

Computer Nodes

对数据进行计算、转换、重塑。

• filter，FilterNodeOptions，移除不符合过滤表达式的行
• project，ProjectNodeOptions，通过执行计算表达式创建新的列，也可以移除、排序列
• aggregate，AggregateNodeOptions，计算整个输入流或数据组的汇总统计信息
• pivot_longer，PivotLongerNodeOptions，通过将某些列转换为附加行来重塑数据

Arrangement Nodes

重新排序、组合或切片数据流。

• hash_john，HashJoinNodeOptions，基于相同的列合并两个输入
• asofjoin，AsofJoinNodeOptions，工具相同的有序列把多个输入合并到第一个输入
• union，合并两个具有相同 Schema 的输入
• order_by，OrderByNodeOptions，对流排序
• fetch，FetchNodeOptions，从流中切片出一系列行

Sink Nodes

Sink 节点会结束一个 Plan。通常不需要创建 Sink Node，会基于 DeclarationToXyz 方法选择。

Conceptions

• Core
  • Buffer
  • Memory Pool
  • Array 数组
  • Scalar 标量，表示有特定类型的单值
  • Record Batch
  • Table
  • Chunked Array
  • Table
• Compute
  • Function
  • Function Registry，注册自定义处理函数
  • Datum
  • Expression
  • ExecBatch
  • BufferSpan
  • ArraySpan
  • ExecSpan
• Acero
  • Declaration
  • ExecPlan
  • ExecNode
• Other

PushGenerator

JoinType

KernelState

KernelContext

KernelInitArgs

ScalarKernel

使用

自定义 Function

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

// Function 帮助类
struct FunctionDoc {
  std::string summary;
  std::string description;
  std::vector<std::string> arg_names;
  std::string options_class; // 参数类
  bool options_required;     // 是否需要参数（Options）
}

// 参数数量标记类
struct Arity {
	Arity Nullary(); // 无参数
  Arity Unary();   // 一个参数
  Arity Binary();  // 两个参数
  Arity Ternary(); // 三个参数
  Arity VarArgs(); // 动态参数
}

// 函数类型
enum Kind {
  SCALAR,
  VECTOR,
  SCALAR_AGGREGATE,
  HASH_AGGREGATE,
  META
}

// 函数基类
Function(std::string name, Function::Kind kind, const Arity& arity, FunctionDoc doc, 
         const FunctionOptions* default_options);
// ScalarFunction 实现了 FunctionImpl<ScalarKernel>
ScalarFunction(std::string name, const Arity& arity, FunctionDoc doc,
               const FunctionOptions* default_options = NULLPTR, bool is_pure = true);

// Kernel
using ArrayKernelExec = Status (*)(KernelContext*, const ExecSpan&, ExecResult*);

TODO

• 需要对 Kernel 做详细了解

附录

• https://cloud.tencent.com/developer/article/2322102

nn

Module

所有神经网络模块的基类，自定义的模型也应把该类作为基类。

1

forward(*input) : 定义每次调用时的计算逻辑，所有子类都应该重写

Sequential

提供了一种简单的方式来按顺序堆叠神经网络的层，用于快速构建一个顺序的神经网络模型。在模型进行前向传播时，nn.Sequential会按照层的顺序依次调用每个层的forward方法，将前一层的输出作为下一层的输入，直到最后一层输出结果。

Linear

ReLU

DDP

• rank，全局进程序号
• local_rank，本机进程序号，可以使用 local_rank = torch.distributed.get_rank() % torch.cuda.device_count() 计算（在每个机器进程数一致时）
• world size，全局并行数

附录

• https://www.alibabacloud.com/blog/alluxio-deep-learning-practices---1-running-pytorch-framework-on-hdfs_596532

• https://medium.com/@binfan_alluxio/deep-learning-at-alibaba-cloud-with-alluxio-running-pytorch-on-hdfs-3a99d7655fd4

• https://zhuanlan.zhihu.com/p/76638962

• https://www.paddlepaddle.org.cn/documentation/docs/zh/api_guides/low_level/layers/sparse_update.html

打包

1. 先编译 C++ 库，生成 metaspore.so
2. 再使用 setuptools 和 wheel 工具，打包 python 库

编译

镜像

• 环境变量

  1 2	export REPOSITORY={hub-repo} export VERSION={version}

• 构建 Dev 镜像（基础环境）

1
2
3

DOCKER_BUILDKIT=1 docker build --network host --build-arg http_proxy=${http_proxy} --build-arg https_proxy=${https_proxy} --build-arg RUNTIME=gpu -f docker/ubuntu20.04/Dockerfile_dev -t $REPOSITORY/metaspore-dev-gpu:${VERSION} .

DOCKER_BUILDKIT=1 docker build --network host --build-arg http_proxy=${http_proxy} --build-arg https_proxy=${https_proxy} --build-arg RUNTIME=cpu -f docker/ubuntu20.04/Dockerfile_dev -t $REPOSITORY/metaspore-dev-cpu:${VERSION} .

• Serving

  • Build 镜像（基于 Dev 镜像进行编译）
  • Service 镜像

• Training

  • Build 镜像（基于 Dev 镜像进行编译）

    1	DOCKER_BUILDKIT=1 docker build --network host --build-arg http_proxy=${http_proxy} --build-arg https_proxy=${https_proxy} -f docker/ubuntu20.04/Dockerfile_training_build --build-arg DEV_IMAGE=$REPOSITORY/metaspore-dev-cpu:${VERSION} -t $REPOSITORY/metaspore-training-build:${VERSION} .

  • Spark Training 镜像

    1

    DOCKER_BUILDKIT=1 docker build --network host --build-arg http_proxy=${http_proxy} --build-arg https_proxy=${https_proxy} -f docker/ubuntu20.04/Dockerfile_training_release --build-arg METASPORE_RELEASE=build --build-arg METASPORE_BUILD_IMAGE=$REPOSITORY/metaspore-training-build:${VERSION} -t $REPOSITORY/metaspore-training-release:${VERSION} --target release .

• Jupyter

  1

  DOCKER_BUILDKIT=1 docker build --network host --build-arg http_proxy=${http_proxy} --build-arg https_proxy=${https_proxy} -f docker/ubuntu20.04/Dockerfile_jupyter --build-arg RELEASE_IMAGE=$REPOSITORY/metaspore-training-release:${VERSION} -t $REPOSITORY/metaspore-training-jupyter:${VERSION} docker/ubuntu20.04

MetaSpore C++

MetaSpore C++ 包含几个模块。

• common
• serving
• metaspore (shared)

common

• globals

  • 定义 gflags 变量

• hashmap

• arrow

• features

metaspore

1. 提供离线训练、在线 serving 的共用代码
2. 离线使用
   1. 使用 pybind11 库定义并绑定 C++ 代码接口
   2. python 代码加载共享库，像调用 python 代码一样调用使用 pybind11 定义的 C++ 接口

Getting Started

文档链接

步骤

• 定义模型（PyTorch Module）
• 定义 Estimator
  • PyTorchEstimator，封装 PyTorch 模型，并在分布式环境下训练（调用 fit 方法并传入 DataFrame 进行训练）
  • 调用 launcher.launch() 在个节点启动 PS 进程（server、worker、coordinator）

定义模型

1
2
3
4
5
6
7
8

embedding_size      : 每个特征组的 embedding size
# MetaSpore 相关
sparse              : ms.EmbeddingSumConcat
sparse.updater      : ms.FTRLTensorUpdater
sparse.initializer  : ms.NormalTensorInitializer
dense.normalization : ms.nn.Normalization
# Torch 相关
dense               : torch.nn.Sequential

初始化内容。

• EmbeddingSumConcat
  • SparseFeatureExtractor
    • 解析原始特征列配置文件
    • 向计算图中添加计算 Hash 特征的 Node
  • EmbeddingBagModule
• TensorUpdater，Sparse & Dense 数据更新类
  • FTRLTensorUpdater
• TensorInitializer，张量初始化器
  • NormalTensorInitializer，归一化张量初始化器
• Normalization，归一化

训练模型

1

PyTorchEstimator

• 定义 PyTorchEstimator
  • module
  • worker / server 数量
  • 模型输出路径
  • Label 列索引

1
2
3
4
5

PyTorchAgent
PyTorchLauncher
PyTorchHelperMixin
PyTorchModel
PyTorchEstimator

核心概念

• JobRunner
• PyTorchEstimator
  • pyspark.ml.base.Estimator
• Launcher
  • PSLauncher
• Agent
• Module
  • EmbeddingOperator
  • TensorUpdater
  • TensorInitializer
  • Normalization
• PyTorchModel
  • pyspark.ml.base.Model
• Metric

C++

common

特征计算相关

FeatureComputeContext 和 FeatureComputeExecContext 都是在 cpp 文件中进行定义。

FeatureComputeContext

FeatureComputeExecContext

FeatureComputeExec

metaspore

Python

embedding

EmbeddingBagModule

1

EmbeddingBagModule -> torch.nn.Module

EmbeddingSumConcat

1

EmbeddingSumConcat -> EmbeddingOperator -> torch.nn.Module

nn

Normalization

others

FTRLTensorUpdater

NormalTensorInitializer

附录

boost

开源方案

• Prophecis
• MetaSpore
• TensorNet
  • 参考一
• TFRA

训练框架

• Tensorfow
• PyTorch
• PaddlePaddle
• MegEngine
• Keras
• MXNet
• CNTK

Inference / Serving

• triton
• TF Serving

MLOps

• cube-studio
• ML-Model-CI

云商产品

• TI-ONE

文章

• 美团 - 一站式机器学习平台建设实践
• TensorFlow在美团外卖推荐场景的GPU训练优化实践
• Distributed hash

训推一体方案

稀疏数据集

• https://www.paddlepaddle.org.cn/documentation/docs/zh/api_guides/low_level/layers/sparse_update.html

文章

• MLOps（六）-回顾2023年开源的MLOps产品、框架、工具与格局变化

神经网络与感知机学习规则

基于神经元模型，提出了感知机学习规则。感知机规则提出了一个可以自动学习的权重优化算法。

感知机算法步骤如下。

1. 初始化权重和偏置项为 0 或很小的随机数
2. 遍历每个训练样本
   1. 计算感知机输出值
   2. 更新权重和偏置项

需要注意的是。

1. 只有当训练数据线性可分时，才能保证感知机具有收敛性
   1. 此时需要设置训练数据集的最大循环次数，或容错次数的阈值，来结束训练
2. 权重、偏置项使用很小的初始化值替代 0 ，如果全是 0 则学习率会失去对决策边界的影响
   1. 学习率只影响权重向量的大小，不影响其方向

错误处理

运行时 API 错误码

调用 CUDA 运行时 API 时，接口返回错误码。

1

__host__ __device__ cudaError_t cudaGetDeviceCount ( int* count ); // 获取设备数量, 返回错误码

错误检查

1
2

__host__ __device__ const char* cudaGetErrorName ( cudaError_t error );     // 获取错误码的枚举名称
__host__ __device__ const char*	cudaGetErrorString ( cudaError_t error );   // 获取错误码的解释描述

定义错误检查函数

1
2
3
4
5
6
7
8
9
10
11
12
13
14

__host__ void error_check_entry() {
  int device_id_in_use;
  error_check(cudaGetDevice(&device_id_in_use), __FILE__, __LINE__);
  error_check(cudaSetDevice(999), __FILE__, __LINE__);
  //  char *p_c;
  //  error_check(cudaMalloc(&p_c, 100), __FILE__, __LINE__);

  cudaDeviceSynchronize();
} /** output
error_check, ok
CUDA error:
        code=101, name=cudaErrorInvalidDevice, description=invalid device ordinal,
        file=/data/code/cook-cuda/src/sample/hello_world.cu, line=51
*/

核函数中的异常

核函数的返回值必须是 void。

1

__host__ __device__ cudaError_t cudaGetLastError ( void ); // 返回最后一次错误码

1
2
3
4
5
6
7
8
9

__global__ void kernel_error_entry() {
  dim3 block(2048);
  print_build_in_vars<<<2, block>>>();  // block size 最大 1024
  error_check(cudaGetLastError(), __FILE__, __LINE__);
} /** output
CUDA error:
        code=9, name=cudaErrorInvalidConfiguration, description=invalid configuration argument,
        file=/data/code/cook-cuda/src/sample/hello_world.cu, line=67
*/

性能评估

事件计时

1
2
3
4
5

__host__ cudaError_t cudaEventCreate ( cudaEvent_t* event );
__host__ __device__ cudaError_t 	cudaEventRecord ( cudaEvent_t event, cudaStream_t stream = 0 );
__host__ cudaError_t cudaEventSynchronize ( cudaEvent_t event );
__host__ cudaError_t cudaEventElapsedTime ( float* ms, cudaEvent_t start, cudaEvent_t end );
__host__ __device__ cudaError_t 	cudaEventDestroy ( cudaEvent_t event );

示例。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

cudaEvent_t start, end;
error_check(cudaEventCreate(&start), __FILE__, __LINE__);
error_check(cudaEventCreate(&end), __FILE__, __LINE__);
error_check(cudaEventRecord(start), __FILE__, __LINE__);
cudaEventQuery(start);

// run GPU Task

error_check(cudaEventRecord(end), __FILE__, __LINE__);
error_check(cudaEventSynchronize(end), __FILE__, __LINE__);
float elapsed_time_ms;
ERROR_CHECK(cudaEventElapsedTime(&elapsed_time_ms, start, end));

printf("elapsed time: %f ms\n", elapsed_time_ms);
ERROR_CHECK(cudaEventDestroy(start));
ERROR_CHECK(cudaEventDestroy(end));

error_check。

1
2
3
4
5
6
7
8

__host__ __device__ cudaError_t error_check(cudaError_t err, const char *fn, int line) {
  if (err != cudaSuccess) {
    printf("CUDA error:\n\tcode=%d, name=%s, description=%s, \n\tfile=%s, line=%d\n", err, cudaGetErrorName(err),
           cudaGetErrorString(err), fn, line);
  }
  return err;
}
#define ERROR_CHECK(exp) error_check(exp, __FILE__, __LINE__)

nvprof

nvprof 是评估 cuda 程序性能的工具。不过目前已经是过时的工具，不适用 compute capability >= 8.0 的设备。新设备适用 nsys 替代。

1

$ nvprof {cuda-program}

nsys

1
2

$ nsys profile {cuda-program} # 运行并记录程序的 profile 到 nsys-rep 文件
$ nsys analyze {nsys-rep}     # 分析 profile 文件

获取 GPU 信息

运行时 API

1

__host__ cudaError_t cudaGetDeviceProperties ( cudaDeviceProp* prop, int  device )

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

__host__ void PrintDeviceInfo() {
  int deviceCount;
  cudaGetDeviceCount(&deviceCount);
  std::cout << "GPU device count: " << deviceCount << std::endl;

  for (int i = 0; i < deviceCount; ++i) {
    // sm: 流式多处理器, Streaming Multiprocessor
    cudaDeviceProp dp{};
    cudaGetDeviceProperties(&dp, i);
    std::cout << "device.0  " << std::endl;
    std::cout << "  sm count: \t\t\t\t" << dp.multiProcessorCount << std::endl;
    std::cout << "  shared memory per block: \t\t" << dp.sharedMemPerBlock / 1024 << "KB" << std::endl;
    std::cout << "  max threads per block:\t\t" << dp.maxThreadsPerBlock << std::endl;
    std::cout << "  max threads per multi processor:\t" << dp.maxThreadsPerMultiProcessor << std::endl;
    std::cout << "  max threads per sm:\t\t\t" << dp.maxThreadsPerMultiProcessor / 32 << std::endl;
    std::cout << "  max blocks per multi processor:\t" << dp.maxBlocksPerMultiProcessor << std::endl;
  }
}