MindSpore自定义算子中的张量维度问题

DechinPhy

发布于 2024-03-13 08:36:50

630

发布于 2024-03-13 08:36:50

文章被收录于专栏：Dechin的专栏Dechin的专栏

技术背景

在前面的几篇博客中，我们介绍了MindSpore框架下使用CUDA来定义本地算子的基本方法，以及配合反向传播函数的使用，这里主要探讨一下MindSpore框架对于CUDA本地算子的输入输出的规范化形式。

测试思路

MindSpore使用的CUDA算子规范化接口形式为：

extern "C" int CustomOps(int nparam, void **params, int *ndims, int64_t **shapes, const char **dtypes,
                         void *stream, void *extra)

也就是说，我们在一个.cu文件中按照这种形式写好函数接口，其中主要是规范化输入输出的形式，然后再将各项输入传给写好的CUDA Kernel函数进行计算并获得返回值。

我们可以使用一个Kernel打印函数的测试案例来说明MindSpore对于输入输出的处理：

#include <iostream>
#define THREADS 1024

__global__ void OpsKernel(const int shape0, const int *input){
    auto i = blockIdx.x * THREADS + threadIdx.x;
    if (i < shape0){
        printf("%d\n", input[i]);
    }
}

在这个函数体内，会把指定大小范围内的input的内容打印出来。

常数输入

首先我们来看一下最简单的常数输入，可以用一个最简单的整数来测试，对应的CUDA算子代码为：

// nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu

#include <iostream>
#define THREADS 1024

__global__ void OpsKernel(const int shape0, const int *input){
    auto i = blockIdx.x * THREADS + threadIdx.x;
    if (i < shape0){
        printf("%d\n", input[i]);
    }
}

extern "C" int CustomOps(int nparam, void **params, int *ndims, int64_t **shapes, const char **dtypes,
                         void *stream, void *extra){
    int *input = static_cast<int*>(params[0]);
    OpsKernel<<<1, THREADS>>>(shapes[0][0], input);
    return 0;
}

调用CUDA算子的Python代码为：

import os
import numpy as np
import mindspore as ms
from mindspore import ops, Tensor, context

context.set_context(mode=context.GRAPH_MODE, device_target="GPU")

CURRENT_PATH = os.path.abspath(__file__)
CustomOps = ops.Custom(CURRENT_PATH.replace(".py", ".so:CustomOps"),
                       out_shape=lambda x:x,
                       out_dtype=ms.int32,
                       func_type="aot")
T0 = Tensor([7], ms.int32)
print (T0)
CustomOps(T0)

运行的指令为：

$ nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu && python3 test_shape.py 
[7]
7

需要注意的是，这里只能给MindSpore内置的几种Tensor变量，如果是直接调用CustomOps(7)会报一个段错误。

高维张量输入

这里一维的张量输入我们就不做讨论了，因为跟前面用到的常数输入本质上是一样的形式。这里我们用一个二维的张量来做一个测试，CUDA代码保持不动，只修改Python代码中的输入：

import os
import numpy as np
import mindspore as ms
from mindspore import ops, Tensor, context

context.set_context(mode=context.GRAPH_MODE, device_target="GPU")

CURRENT_PATH = os.path.abspath(__file__)
CustomOps = ops.Custom(CURRENT_PATH.replace(".py", ".so:CustomOps"),
                       out_shape=lambda x:x,
                       out_dtype=ms.int32,
                       func_type="aot")
T0 = Tensor(np.arange(12).reshape((4, 3)), ms.int32)
print (T0)
CustomOps(T0)

运行结果为：

$ nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu && python3 test_shape.py 
[[ 0  1  2]
 [ 3  4  5]
 [ 6  7  8]
 [ 9 10 11]]
0
1
2
3

需要注意的是，我们在CUDA的打印函数中设置的打印输出大小是输入张量的第一个维度的大小，我们给的是一个(4,3)大小的张量，因此会顺序打印4个数出来。这里我们也能够发现MindSpore在进行输入的规范化的时候，会自动压平输入的张量变成一个维度。因此这里的调用代码等价于先对输入张量做一个reshape，然后再把第一个维度对应大小的张量元素打印出来。如果要打印所有的元素也很简单，可以修改一下CUDA代码：

// nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu

#include <iostream>
#define THREADS 1024

__global__ void OpsKernel(const int shape0, const int *input){
    auto i = blockIdx.x * THREADS + threadIdx.x;
    if (i < shape0){
        printf("%d\n", input[i]);
    }
}

extern "C" int CustomOps(int nparam, void **params, int *ndims, int64_t **shapes, const char **dtypes,
                         void *stream, void *extra){
    int *input = static_cast<int*>(params[0]);
    int elements = 1;
    for (int i=0; i<ndims[0]; i++){
        elements *= shapes[0][i];
    }
    OpsKernel<<<1, THREADS>>>(elements, input);
    return 0;
}

通过定义一个elements变量用于存储对应张量的元素数量，然后再逐一打印出来即可，执行结果为：

$ nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu && python3 test_shape.py 
[[ 0  1  2]
 [ 3  4  5]
 [ 6  7  8]
 [ 9 10 11]]
0
1
2
3
4
5
6
7
8
9
10
11

输出规范化

当我们使用ops.Custom算子时，如果指定了out_dtype和out_shape，那么算子会自动帮我们分配好相应的device memory空间。那么我们在CUDA计算的时候可以直接修改对应的内存空间：

// nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu

#include <iostream>
#define THREADS 1024

__global__ void OpsKernel(const int shape0, const int *input, float *output){
    auto i = blockIdx.x * THREADS + threadIdx.x;
    if (i < shape0){
        output[i] = input[i] * 0.5;
    }
}

extern "C" int CustomOps(int nparam, void **params, int *ndims, int64_t **shapes, const char **dtypes,
                         void *stream, void *extra){
    int *input = static_cast<int*>(params[0]);
    float *output = static_cast<float*>(params[1]);
    int elements = 1;
    for (int i=0; i<ndims[0]; i++){
        elements *= shapes[0][i];
    }
    OpsKernel<<<1, THREADS>>>(elements, input, output);
    return 0;
}

这里我们对算子的功能做了一点调整，我们输出的结果是整个张量的元素值乘以0.5，同时也把一个整形变量转化成了一个浮点型变量。其运行Python代码也要做一点调整：

import os
import numpy as np
import mindspore as ms
from mindspore import ops, Tensor, context

context.set_context(mode=context.GRAPH_MODE, device_target="GPU")

CURRENT_PATH = os.path.abspath(__file__)
CustomOps = ops.Custom(CURRENT_PATH.replace(".py", ".so:CustomOps"),
                       out_shape=lambda x:x,
                       out_dtype=ms.float32,
                       func_type="aot")
T0 = Tensor(np.arange(12).reshape((4, 3)), ms.int32)
print (T0)
output = CustomOps(T0)
print (output)

这里主要是修改了out_dtype为浮点型，这里如果写错了，会直接导致内存溢出。上述代码的运行结果为：

$ nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu && python3 test_shape.py 
[[ 0  1  2]
 [ 3  4  5]
 [ 6  7  8]
 [ 9 10 11]]
[[0.  0.5 1. ]
 [1.5 2.  2.5]
 [3.  3.5 4. ]
 [4.5 5.  5.5]]

可以看到这里输出的张量形状是跟输入保持一致的，即时这个输入张量在经过MindSpore的Custom算子接口时已经被压平成一个一维张量，但是因为我们设置了out_shape=lambda x:x，这表示输出的张量shape跟输入的张量shape一致，当然，直接用Python的列表来给out_shape赋值也是可以的。例如我们写一个输入输出不同shape的案例：

// nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu

#include <iostream>
#define THREADS 1024

__global__ void OpsKernel(const int shape0, const int *input, int *output){
    auto i = blockIdx.x * THREADS + threadIdx.x;
    if (i < shape0){
        atomicAdd(&output[0], input[i]);
    }
}

extern "C" int CustomOps(int nparam, void **params, int *ndims, int64_t **shapes, const char **dtypes,
                         void *stream, void *extra){
    int *input = static_cast<int*>(params[0]);
    int *output = static_cast<int*>(params[1]);
    int elements = 1;
    for (int i=0; i<ndims[0]; i++){
        elements *= shapes[0][i];
    }
    OpsKernel<<<1, THREADS>>>(elements, input, output);
    return 0;
}

这个Kernel函数的主要功能是通过一个atomicAdd函数，把输入张量的所有元素做一个求和，这样输出的张量的shape只有1，对应的Python调用形式也要做一定的调整：

import os
import numpy as np
import mindspore as ms
from mindspore import ops, Tensor, context

context.set_context(mode=context.GRAPH_MODE, device_target="GPU")

CURRENT_PATH = os.path.abspath(__file__)
CustomOps = ops.Custom(CURRENT_PATH.replace(".py", ".so:CustomOps"),
                       out_shape=[1],
                       out_dtype=ms.int32,
                       func_type="aot")
T0 = Tensor(np.arange(12).reshape((4, 3)), ms.int32)
print (T0)
output = CustomOps(T0)
print (output)

由于atomicAdd(addr, element)原子操作要求输入输出的类型要一致，因此这里我们还是使用的int类型的output，输出结果如下所示：

$ nvcc --shared -Xcompiler -fPIC -o test_shape.so test_shape.cu && python3 test_shape.py 
[[ 0  1  2]
 [ 3  4  5]
 [ 6  7  8]
 [ 9 10 11]]
[66]

总结概要

当我们使用GPU进行快速运算时，虽然可以用成熟的深度学习框架如MindSpore和PyTorch等进行实现，但其实从速度上来说，最快不过直接使用C/C++的CUDA来实现。也正是因为如此，在MindSpore框架中支持了对CUDA实现的算子的直接调用，只是在格式规范上有一定的要求。本文主要介绍MindSpore调用本地CUDA算子的一些规范化和技巧。