CUDA samples系列 0.4 cppOverload

基础知识
- int与int2类型
- 强制类型转换
- 函数指针
- 函数重载
- 设定核函数属性
- ShareMemory
源代码解析

基础知识

这份代码介绍了核函数重载的方法，先介绍一些需要的基础知识。

int与int2类型

int是四个字节，32位；

int2类型是2个int，可以分为2个int：

int2 position；
position.x = 1;
position.y = 3;

强制类型转换

reinterpret_cast为强制类型转换符，他只是把指针以新的方式进行解析，并不会做更多的动作。

int *hInput  = NULL;
cudaMallocHost(&hInput , sizeof(int)*N*2);
check_func2(reinterpret_cast<int2 *>(hInput), hOutput, a)；

上面这里只是把hInput以int2进行解析，原来*（hInput+1）向前前进一个int，现在向前前进2个int。这段代码里，最终传进check_func2()的第一个参数，就是 int2*类型了。

函数指针

函数指针的用法为：

int func(int x); /* 声明一个函数，记得定义 */
int (*f) (int x); /* 声明一个函数指针 */
f=func; /* 将func函数的首地址赋给指针f */
(*f) （３）；／／(*f)取地址解析出这个函数，然后直接调用这个函数

一段简单的代码如下：

#include <stdio.h>
#include <string.h>
 char * fun(char * p1,char * p2)
{
　　int i = 0;
　　i = strcmp(p1,p2);
 
　　if (0 == i)
　　{
　　　　return p1;
　　}
　　else
　　{
　　　　return p2;
　　}
}
 
int main()
{
　　char * (*pf)(char * p1,char * p2);
　　pf = &fun;　
　　/*或者pf＝fun；一样的效果；这是因为函数名被编译之后其实就是一个地址，所以这里两种用法没有本质的差别。*/
　　(*pf) ("aa","bb");
　　return 0;
}

函数重载

写几个相同函数名，但是不同函数参数的函数就是了。

源代码中有这么几个simple_kernel函数，都是核函数（在GPU上运行的函数）：

__global__ void simple_kernel(const int *pIn, int *pOut, int a)；
__global__ void simple_kernel(const int2 *pIn, int *pOut, int a)；
__global__ void simple_kernel(const int *pIn1, const int *pIn2, int *pOut, int a)；

与之配合，定义了3个函数指针：

void (*func1)(const int *, int *, int);
    void (*func2)(const int2 *, int *, int);
    void (*func3)(const int *, const int *, int *, int);

设定核函数属性

用到了一个函数，设定核函数运行时，更多的使用什么内存

官方文档介绍了这个函数。

CUDA samples系列 0.4 cppOverload基础知识源代码解析

获取函数属性：

memset(&attr, 0, sizeof(attr));
    checkCudaErrors(cudaFuncGetAttributes(&attr, *func2));
    OUTPUT_ATTR(attr);

输出显示属性的函数如下：

#define OUTPUT_ATTR(attr)  \
    printf("Shared Size:   %d\n", (int)attr.sharedSizeBytes);   \
    printf("Constant Size: %d\n", (int)attr.constSizeBytes);                 \
    printf("Local Size:    %d\n", (int)attr.localSizeBytes);                 \
    printf("Max Threads Per Block: %d\n", attr.maxThreadsPerBlock);          \
    printf("Number of Registers: %d\n", attr.numRegs);                       \
    printf("PTX Version: %d\n", attr.ptxVersion);                            \
    printf("Binary Version: %d\n", attr.binaryVersion);                      \

ShareMemory

之前，每每说到，拷贝到GPU内存空间，都指的是全局内存；

现在介绍下共享内存ShareMemory，他的空间非常小，每个block有一块几十kB的空间，但是速度非常快；不可以用cudaMemcpy来直接复制数据过去，他仅仅是在计算时作为一个储物台存在的。他最大的特点是：可以被同一块block里的任意进程访问到。

CUDA samples系列 0.4 cppOverload基础知识源代码解析

这是GPU的硬件逻辑图，其中的SM就是一块block，shared memory非常小，非常快，而且同一个SM里的所有的sp都可以访问到。

比如下面这个核函数：

__global__ void simple_kernel(const int *pIn, int *pOut, int a)
{
    __shared__ int sData[THREAD_N];
    int tid = threadIdx.x + blockDim.x*blockIdx.x;

    sData[threadIdx.x] = pIn[tid];
    __syncthreads();

    pOut[tid] = sData[threadIdx.x]*a + tid;;
}

这是使用共享内存的方法，直接在核函数里定义一个大小为THREAD_N个int的空间，THREAD_N是设定的一个block有多少个线程，源代码中为

#define THREAD_N 256

。

源代码解析

首先是“cppOverload.cu”：

#define THREAD_N 256 //每个block里设定256个线程
#define N 1024	//总共计算1024个数
#define DIV_UP(a, b) (((a) + (b) - 1) / (b))//定义函数，这是标准的计算block个数的函数

	//以下引用一段cuda系列第三篇的讲解：
	/*设定block数，线程数；保证总线程数大于矢量的长度，这样每个矢量中的元素都会被计算到
	每个线程计算矢量中的一个元素
	这里设定线程与block数是一个常用的方法：
	先设定threadsPerBlock为某个值，然后计算blocksPerGrid，公式为：
	blocksPerGrid = （要计算的矢量长度+threadsPerBlock-1）/threadsPerBlock
	这样可以保证总共可调用线程数>=要计算的矢量长度 */


// Includes, system
#include <stdio.h>
#include <helper_cuda.h>
#include <helper_string.h>
#include <helper_math.h>
#include "cppOverload_kernel.cuh"

const char *sampleName = "C++ Function Overloading";

#define OUTPUT_ATTR(attr)  \
    printf("Shared Size:   %d\n", (int)attr.sharedSizeBytes);   \
    printf("Constant Size: %d\n", (int)attr.constSizeBytes);                 \
    printf("Local Size:    %d\n", (int)attr.localSizeBytes);                 \
    printf("Max Threads Per Block: %d\n", attr.maxThreadsPerBlock);          \
    printf("Number of Registers: %d\n", attr.numRegs);                       \
    printf("PTX Version: %d\n", attr.ptxVersion);                            \
    printf("Binary Version: %d\n", attr.binaryVersion);                      \
     

//三个cpu检查函数，用来检查gpu核函数的运算结果是否足够准确
bool check_func1(int *hInput, int *hOutput, int a)
{
    for (int i = 0; i < N; ++i)
    {
        int cpuRes = hInput[i]*a + i;

        if (hOutput[i] != cpuRes)
        {
            return false;
        }
    }

    return true;
}

bool check_func2(int2 *hInput, int *hOutput, int a)
{
    for (int i = 0; i < N; i++)
    {
        int cpuRes = (hInput[i].x + hInput[i].y)*a + i;

        if (hOutput[i] != cpuRes)
        {
            return false;
        }
    }

    return true;
}

bool check_func3(int *hInput1, int *hInput2, int *hOutput, int a)
{
    for (int i = 0; i < N; i++)
    {
        if (hOutput[i] != (hInput1[i] + hInput2[i])*a + i)
        {
            return false;
        }
    }

    return true;
}

int main(int argc, const char *argv[])
{
    int *hInput  = NULL;
    int *hOutput = NULL;
    int *dInput  = NULL;
    int *dOutput = NULL;

    printf("%s starting...\n", sampleName);

	//获取可用的GPU设备信息
    int deviceCount;
    checkCudaErrors(cudaGetDeviceCount(&deviceCount));
    printf("DevicecheckCudaErrors Count: %d\n", deviceCount);

    int deviceID = findCudaDevice(argc, argv);
	cudaDeviceProp prop;
	checkCudaErrors(cudaGetDeviceProperties(&prop, deviceID));
	if (prop.major < 2)    
    {
        printf("ERROR: cppOverload requires GPU devices with compute SM 2.0 or higher.\n");
        printf("Current GPU device has compute SM%d.%d, Exiting...", prop.major, prop.minor);
        exit(EXIT_WAIVED);
    }
	
    checkCudaErrors(cudaSetDevice(deviceID));

	// 分配内存空间
    // Allocate device memory
    checkCudaErrors(cudaMalloc(&dInput , sizeof(int)*N*2));
    checkCudaErrors(cudaMalloc(&dOutput, sizeof(int)*N));

    // Allocate host memory
    checkCudaErrors(cudaMallocHost(&hInput , sizeof(int)*N*2));
    checkCudaErrors(cudaMallocHost(&hOutput, sizeof(int)*N));

	// 设定数据的数值1-2048
    for (int i = 0; i < N*2; i++)
    {
        hInput[i] = i;
    }
	
    // Copy data from host to device
    checkCudaErrors(cudaMemcpy(dInput, hInput, sizeof(int)*N*2, cudaMemcpyHostToDevice));

    // Test C++ overloading
    bool testResult = true;
    bool funcResult = true;
    int a = 1;
	
	// 3个函数指针
    void (*func1)(const int *, int *, int);
    void (*func2)(const int2 *, int *, int);
    void (*func3)(const int *, const int *, int *, int);
    struct cudaFuncAttributes attr;//记录函数属性的结构体

    // 重载第一个函数
    func1 = simple_kernel;//由于func的参数类型与第一个simple_kernel吻合，这里将会与之绑定
    memset(&attr, 0, sizeof(attr)); 
	//CacheConfig设定为优先使用共享内存
    checkCudaErrors(cudaFuncSetCacheConfig(*func1, cudaFuncCachePreferShared));
	//获取函数的属性
    checkCudaErrors(cudaFuncGetAttributes(&attr, *func1));
    OUTPUT_ATTR(attr);//输出属性信息
    (*func1)<<<DIV_UP(N, THREAD_N), THREAD_N>>>(dInput, dOutput, a);
    checkCudaErrors(cudaMemcpy(hOutput, dOutput, sizeof(int)*N, cudaMemcpyDeviceToHost));
    funcResult = check_func1(hInput, hOutput, a);//cpu检查一下gpu计算的是否够准确
    printf("simple_kernel(const int *pIn, int *pOut, int a) %s\n\n", funcResult ? "PASSED" : "FAILED");
    testResult &= funcResult;

    // overload function 2
    func2 = simple_kernel;
    memset(&attr, 0, sizeof(attr));
    checkCudaErrors(cudaFuncSetCacheConfig(*func2, cudaFuncCachePreferShared));
    checkCudaErrors(cudaFuncGetAttributes(&attr, *func2));
    OUTPUT_ATTR(attr);
    (*func2)<<<DIV_UP(N, THREAD_N), THREAD_N>>>((int2 *)dInput, dOutput, a);
    checkCudaErrors(cudaMemcpy(hOutput, dOutput, sizeof(int)*N, cudaMemcpyDeviceToHost));
    funcResult = check_func2(reinterpret_cast<int2 *>(hInput), hOutput, a);
    printf("simple_kernel(const int2 *pIn, int *pOut, int a) %s\n\n", funcResult ? "PASSED" : "FAILED");
    testResult &= funcResult;

    // overload function 3
    func3 = simple_kernel;
    memset(&attr, 0, sizeof(attr));
    checkCudaErrors(cudaFuncSetCacheConfig(*func3, cudaFuncCachePreferShared));
    checkCudaErrors(cudaFuncGetAttributes(&attr, *func3));
    OUTPUT_ATTR(attr);
    (*func3)<<<DIV_UP(N, THREAD_N), THREAD_N>>>(dInput, dInput+N, dOutput, a);
    checkCudaErrors(cudaMemcpy(hOutput, dOutput, sizeof(int)*N, cudaMemcpyDeviceToHost));
    funcResult = check_func3(&hInput[0], &hInput[N], hOutput, a);
    printf("simple_kernel(const int *pIn1, const int *pIn2, int *pOut, int a) %s\n\n", funcResult ? "PASSED" : "FAILED");
    testResult &= funcResult;

    checkCudaErrors(cudaFree(dInput));
    checkCudaErrors(cudaFree(dOutput));
    checkCudaErrors(cudaFreeHost(hOutput));
    checkCudaErrors(cudaFreeHost(hInput));

    checkCudaErrors(cudaDeviceSynchronize());

    exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
}

还有个“cppOverload_kernel.cuh”：

__global__ void simple_kernel(const int *pIn, int *pOut, int a)
{
    __shared__ int sData[THREAD_N];
    int tid = threadIdx.x + blockDim.x*blockIdx.x;

    sData[threadIdx.x] = pIn[tid];
    __syncthreads();

    pOut[tid] = sData[threadIdx.x]*a + tid;;
}

__global__ void simple_kernel(const int2 *pIn, int *pOut, int a)
{
    __shared__ int2 sData[THREAD_N];
    int tid = threadIdx.x + blockDim.x*blockIdx.x;

    sData[threadIdx.x] = pIn[tid];
    __syncthreads();

    pOut[tid] = (sData[threadIdx.x].x + sData[threadIdx.x].y)*a + tid;;
}

__global__ void simple_kernel(const int *pIn1, const int *pIn2, int *pOut, int a)
{
    __shared__ int sData1[THREAD_N];
    __shared__ int sData2[THREAD_N];
    int tid = threadIdx.x + blockDim.x*blockIdx.x;

    sData1[threadIdx.x] = pIn1[tid];
    sData2[threadIdx.x] = pIn2[tid];
    __syncthreads();

    pOut[tid] = (sData1[threadIdx.x] + sData2[threadIdx.x])*a + tid;
}

CUDA samples系列 0.4 cppOverload基础知识源代码解析

CUDA samples系列 0.4 cppOverload

基础知识

int与int2类型

强制类型转换

函数指针

函数重载

设定核函数属性

ShareMemory

源代码解析

继续阅读

Ubuntu下安装cuda历程记录

CUDA学习笔记(5) 原子操作

DAT（NIPS 2018）视频目标跟踪源码运行笔记1. 论文基本信息2. 运行环境介绍3. 准备4. 配置5. 运行6. 可能出现的问题及解决方法

Ubuntu16.04+Pytorch1.4.0+cuda10.0的pip安装

【CUDA-C/C++】任意维度矩阵乘

WSL2 的docker里使用显卡的安装

关于安装pytorch的一些问题总结

jetson nano ubuntu 安装opencv4 cuda10 pytorch

Jetson Nano Ubuntu编译OpenCV4.4.0+opencv_contrib(带CUDA)

本科、硕士、博士的区别是什么？

yolov7 tensorrt模型加速部署【实战】

linxu下CUDA静态库-上

《cuda并行程序设计》勘误（3）

《cuda并行程序设计》勘误（2）

ImportError: libcublas.so.9.0: cannot open shared object file: No such file or directory（完美解决）

一种解决思路： ImportError: libcublas.so.10.0: cannot open shared object file: No such file