CUDA学习(十三):一步一步为CUDA提速


笔者CUDA学习系列汇总传送门(持续更新):编程语言|CUDA入门


笔者在CUDA学习(十):向量内积的多种方法实现
CUDA(六):多种方法实现向量加
两篇文章中,已经用了多种方法来实现向量内积和向量加,但一直没有详细记录各种方法的区别以及这些方法之间的速度差距等等.
这里笔者再做一份笔记,浅谈cuda加速的基本技巧.
并记录下各种方法快慢的原理。


一、前言

  • 本文通过循环1000次来统计运行时间,且不考虑内存搬运的时间,所采用的时间计算函数为
long get_time()
{
    struct timeval t_start;
    gettimeofday(&t_start, NULL);

    long end = ((long)t_start.tv_sec) * 1000 + (long)t_start.tv_usec / 1000;;
    return end;
}
  • 由于笔者的方向是图像处理,所以本文采用图像为矩阵的数据来源,具体步骤如下:

1、opencv读取图像,并resize到(512, 512)
2、BGR转GRAY,再转float
3、除以255

 cv::Mat matBgrImg = cv::imread("./DEMO.jpg");
    
 int nStep = matBgrImg.step;
printf("matF32 h = %d, w = %d, channel = %d, step = %d \n", matBgrImg.rows, matBgrImg.cols ,\
                     matBgrImg.channels(),nStep);

 cv::resize(matBgrImg, matBgrImg, cv::Size(512, 512));
 cv::Mat matGrayIMg;
 cv::cvtColor(matBgrImg, matGrayIMg, cv::COLOR_BGR2GRAY);
 cv::Mat matF32;
 matGrayIMg.convertTo( matF32, CV_32FC1);
 matF32 = matF32 / 255.;
  • 本文为了代码更加直观,采用for的循环方式。用for多了一个变量,笔者建议使用while进行循环

二、opencv对图像求和 41ms

直接调用sum函数, 只需要41ms。

fSum = sum(matF32)[0];

三、最愚蠢的方法,直接在GPU内循环求和 31418ms!

这种直接在GPU内循环求和,完全没有发挥GPU多线程的优势,运行速度达到31418ms,是Opencv的1000倍。

  • 为什么这样写GPU这么慢呢?

这由GPU的架构体系决定的,在CUDA中,一般的数据赋值到显卡内存部分,称为 global memory。这些内存是没有cache的,而且存取global memory所需要的时间**(latency)**是非常长的。由于此程序只有一个thread,所以每次读取global memory的内容,就需要等实际读取到数据、累加到Sum后,才能进行下一步。所以,这是一种最愚蠢的方法,运行的时间特别长。

/*================================
* @brief 单线程运算,速度究极慢
*
* @param src 
* @param dst
* @return 
=================================*/
__global__ void vector_add_gpu_2(float *a, float *b, int n)
{
    for(int i = 0; i < n; ++i)
    {
        b[0] += a[i];
    }
}

四、初级加速:单block,多thread

此方法初步将cuda的并行利用起来,在256个线程下,运行速度达到173ms。

  • 在这里,得先了解一下GOU内存的存取模式。

显卡上的内存是DRAM,因此最具有效率的存取方式是以连续的方式存取。
但考虑到thread的执行方式,当一个thread在等待内存数据时,GPU会切换
到下一个thread,也就是说,实际上执行的顺序类似
thread0-> thread1 -> thread2
所以呢,我们需要让thread0 读取第0个数字,thread1 读取第一个数字。这样才能保证数据读取的效率。

/*================================
* @brief 单block,多thread,thread间用原子操作加
*
* @param src 
* @param dst
* @return 
=================================*/
__global__ void  nv_SumOfv2(float *pData, float *pSum, int N)
{
     int tid = threadIdx.x;
    float fSum = 0.0;
    const int nStep = blockDim.x;

    while(tid < N)
    {
        fSum += pData[tid];
        tid += nStep;
    }
    atomicAdd(pSum, fSum); // 原子操作
}

仅仅这么一个操作,运行速度就提高了上千倍,但离opencv求和的标准还有很大的距离

五、中级加速:多block多thread 36ms!

在笔者以前的笔记中,已经提到过block的概念。
我们已经知道,在CUDA中,block是可以分组的,且每一个block中的所有thread,具有一个可以共享share memory,并且进行同步操作。 注意:不同block之间的thread则不行。
笔者这里举了个简单例子来描述block和thread之间并行的关系。
block看成是小分队1,、小分队2.。每个小分队里又有成员1、成员2.(thread)。
一个小分队里的所有人同时开始搬砖,速度肯定比不上多个小分队同时搬砖。
thread完成自己的工作了,可thread搬的砖如何堆在一起呢(求总和)?这时候,cuda中的原子操作来了,都放在那,让我来。注意:cuda中的原子操作是顺序执行的,虽然代码简洁,但很可能会影响程序的运行效率。

/*================================
* @brief 多block,多thread, 
* 利用原子操作 代码比较简单,但是原子操作对数据的访问是串行的,频繁的原子操作会影响性能
* 此时原子操作了 blockDim.x * gridDim.x 次
* @param src 
* @param dst
* @return 
=================================*/
__global__ void sumOfCuda(float *pData, float *pSum, int N)
{
    
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int nStep = blockDim.x * gridDim.x;

    double dSum = 0.0;
    for (int i = tid; i < N; i+=nStep)
    {
        dSum += pData[i];
    }
    atomicAdd(pSum, dSum); // 对所有线程进行原子操作求和
}

此番操作,在16个block,256个thread下,1000次运行下,速度达到 36ms,已经超过opencv求和的速度了。还有加速空间么?当然有,block间也可以同步加速呀

六、高级加速:利用shared memory 多block多thread 21ms!

thread里面多线程归约,归约后将各个thread的结果保存到共享内存里,
然后再对block内归约, 最后将各个block的结果用原子操作求和。
利用__shared__ 声明的变量表示这是 shared memory,是一个block中每个thread都共享的内存。这是GPU上的内存,存取的速度相当快,不需要担心latency的问题__syncthreads()是一个cuda内部的函数,表示block中的所有thread都要同步到这个点,
才能继续执行下一步操作。 在操作共享内存时,还需要注意bank conflict的问题

/*================================
* @brief thread 里面多线程归约,归约后将各个thread的结果保存到共享内存里,
* 然后再对block归约, 最后将各个block的结果用原子操作求和
* 利用__shared__ 声明的变量表示这是 shared memory,是一个block中每个thread
* 都共享的内存, 这会使用GPU上的内存,存取的速度相当快,不需要担心latency的问题
* __syncthreads()是一个cuda内部的函数,表示block中的所有thread都要同步到这个点,
* 才能继续执行下一步操作。 在操作共享内存时,需要注意bank conflict的问题
* @param src 
* @param dst
* @return 
=================================*/
__global__ void sumOfCuda_v2(float *pfData, float *pSum, int N)
{
    // printf("blockDim.x = %d\n", blockDim.x);
    __shared__ double share_dTemp[THREAD_NUM];
    const int nStep = gridDim.x * blockDim.x;
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    double dTempSum = 0.0;
    for (int i = tid; i < N; i+=nStep)
    {
        dTempSum += pfData[i];
    }
    share_dTemp[threadIdx.x] = dTempSum; 
    __syncthreads();// 同步操作,等待上面执行完成
    // 此时每个block内的每一个线程,都放了各自的求和
    // 然后需要对每个block内的线程进行归约
    // 每个block内有 blockDim.x 个线程, 也就是对每个block内的

    for (int i = blockDim.x / 2; i !=0; i/=2)
    {
        if (threadIdx.x < i)
        {
            share_dTemp[threadIdx.x] += share_dTemp[threadIdx.x + i];
        }
        __syncthreads();
    }

    if(0 == threadIdx.x)
    {
         atomicAdd(pSum, share_dTemp[0]);
    }
}

16个block,256个thread下,运行一千次,速度进一步优化,达到21ms。

以上程序还能进一步优化,待笔者对CUDA研究透彻后,再仔细道来。

最后附上demo的运行例子

#include "npp.h"
#include <cuda_runtime.h>
#include "cuda_runtime_api.h"
#include "device_launch_parameters.h"
#include "opencv2/opencv.hpp"
#include <time.h>
#include <sys/time.h>

// int nGpuId = 1;
// cudaSetDevice(nGpuId);

long get_time()
{
    struct timeval t_start;
    gettimeofday(&t_start, NULL);

    long end = ((long)t_start.tv_sec) * 1000 + (long)t_start.tv_usec / 1000;;
    return end;
}
/*================================
* @brief 多block,多thread, 
* 利用原子操作 代码比较简单,但是原子操作对数据的访问是串行的,频繁的原子操作会影响性能
* 此时原子操作了 blockDim.x * gridDim.x 次
* @param src 
* @param dst
* @return 
=================================*/
__global__ void sumOfCuda(float *pData, float *pSum, int N)
{
    
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int nStep = blockDim.x * gridDim.x;

    double dSum = 0.0;
    for (int i = tid; i < N; i+=nStep)
    {
        dSum += pData[i];
    }
    atomicAdd(pSum, dSum); // 对所有线程进行原子操作
}

const int N = 512*512;

const int THREAD_NUM = 256;
const int BLOCK_NUM = 16;

/*================================
* @brief thread 里面多线程归约,归约后将各个thread的结果保存到共享内存里,
* 然后再对block归约, 最后将各个block的结果用原子操作求和
* 利用__shared__ 声明的变量表示这是 shared memory,是一个block中每个thread
* 都共享的内存, 这会使用GPU上的内存,存取的速度相当快,不需要担心latency的问题
* __syncthreads()是一个cuda内部的函数,表示block中的所有thread都要同步到这个点,
* 才能继续执行下一步操作。 在操作共享内存时,需要注意bank conflict的问题
* @param src 
* @param dst
* @return 
=================================*/
__global__ void sumOfCuda_v2(float *pfData, float *pSum, int N)
{
    // printf("blockDim.x = %d\n", blockDim.x);
    __shared__ double share_dTemp[THREAD_NUM];
    const int nStep = gridDim.x * blockDim.x;
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    double dTempSum = 0.0;
    for (int i = tid; i < N; i+=nStep)
    {
        dTempSum += pfData[i];
    }
    share_dTemp[threadIdx.x] = dTempSum; 
    __syncthreads();// 同步操作,等待上面执行完成
    // 此时每个block内的每一个线程,都放了各自的求和
    // 然后需要对每个block内的线程进行归约
    // 每个block内有 blockDim.x 个线程, 也就是对每个block内的

    for (int i = blockDim.x / 2; i !=0; i/=2)
    {
        if (threadIdx.x < i)
        {
            share_dTemp[threadIdx.x] += share_dTemp[threadIdx.x + i];
        }
        __syncthreads();
    }

    if(0 == threadIdx.x)
    {
         atomicAdd(pSum, share_dTemp[0]);
    }
}


void test_v1()
{
    cv::Mat matBgrImg = cv::imread("./demo.jpg");
    
    cv::resize(matBgrImg, matBgrImg, cv::Size(512, 512));
    cv::Mat matGrayIMg;
    cv::cvtColor(matBgrImg, matGrayIMg, cv::COLOR_BGR2GRAY);
    cv::Mat matF32;
    matGrayIMg.convertTo( matF32, CV_32FC1);
    matF32 = matF32 / 255.;
    int nH = matF32.rows;
    int nW = matF32.cols;
    int nStep = matF32.step;
    printf("matF32 h = %d, w = %d, channel = %d, step = %d \n", matF32.rows, matF32.cols ,\
                    matF32.channels(),nStep);

    void *pvData = malloc(1 * 1 * N *sizeof(float));
     memcpy(pvData, (unsigned char *)matF32.data, N * sizeof(float));
    float *pfData_dev = NULL;

    cudaMalloc((void **)&pfData_dev, N * sizeof(float));
    cudaMemcpy(pfData_dev, pvData, N * sizeof(float), cudaMemcpyHostToDevice);

    float fSum = 0.0;
    long time_start = get_time();
    float *pfSum_dev = NULL;
    cudaMalloc((void **)&pfSum_dev, THREAD_NUM * sizeof(float));

    printf("block num = %d, thread num = %d \n", BLOCK_NUM, THREAD_NUM);

    for (int i = 0;  i < 1000; ++i)
    {
        cudaMemset(pfSum_dev, 0, THREAD_NUM * sizeof(float));
        // printf("begine %d, %d\n", THREAD_NUM, THREAD_NUM);

        sumOfCuda<<<BLOCK_NUM, THREAD_NUM>>> (pfData_dev, pfSum_dev, N);
        // std::cout << "w" << std::endl;
        // vector_add_gpu_2<<<1, 1>>> (pfData_dev, pfSum_dev, N);

        float pfSum[THREAD_NUM] = {0.0};
        cudaMemcpy(pfSum, pfSum_dev, THREAD_NUM * sizeof(float), cudaMemcpyDeviceToHost);
        fSum = 0.0;
        for (int j = 0; j < THREAD_NUM; ++j)
        {
            fSum += pfSum[j];
        }
    }
    printf("sum = %f\n", fSum);
     std::cout << "costime is " << get_time() - time_start << std::endl;
    printf("sum = %f\n", fSum);
    return ;
}

int main()
{
    test_v1();
    return 0;
}
  • 6
    点赞
  • 49
    收藏
    觉得还不错? 一键收藏
  • 6
    评论
### 回答1: 实现cv::seamlessClone可以使用OpenCV库中提供的CUDA函数进行实现。 以下是一个简单的示例代码: ``` #include <opencv2/opencv.hpp> #include <opencv2/cudaimgproc.hpp> #include <opencv2/cudaarithm.hpp> int main(int argc, char** argv) { cv::Mat src = cv::imread("src.jpg"); cv::Mat dst = cv::imread("dst.jpg"); cv::Mat mask = cv::imread("mask.jpg", 0); cv::cuda::GpuMat src_gpu, dst_gpu, mask_gpu, result_gpu; src_gpu.upload(src); dst_gpu.upload(dst); mask_gpu.upload(mask); cv::cuda::seamlessClone(src_gpu, dst_gpu, mask_gpu, cv::Point(dst.cols / 2, dst.rows / 2), result_gpu, cv::cuda::NORMAL_CLONE); cv::Mat result_cpu; result_gpu.download(result_cpu); cv::imshow("result", result_cpu); cv::waitKey(0); return 0; } ``` 在此代码中,我们首先加载了原始图像、目标图像和掩码图像,然后将它们上传到GPU。接下来,我们调用`cv::cuda::seamlessClone`函数,并将结果下载到CPU上的矩阵中。最后,我们使用`cv::imshow`函数显示结果。 ### 回答2: 使用CUDA代码实现cv::seamlessClone需要以下步骤: 1. 首先,将输入图像和目标图像从主机内存复制到CUDA设备内存中。可以使用cudaMemcpy函数进行内存拷贝。 2. 在CUDA设备上创建一个输出图像的内存空间,并使用cudaMalloc函数为其分配内存。 3. 将输入图像和目标图像的像素数据分别传送到CUDA设备内存中。可以使用cudaMemcpy2D函数将二维图像数据传送到设备。 4. 在CUDA设备上创建一个内核函数,用来计算图像中的每个像素点的融合颜色。该函数可以根据融合算法的不同,使用不同的插值方法来计算像素点的新颜色。 5. 调用内核函数,对每个像素点进行并行计算,计算结果存储在输出图像内存中。 6. 最后,将输出图像的像素数据从设备内存复制到主机内存中。可以使用cudaMemcpy2D函数将二维图像数据从设备复制到主机内存。 7. 在主机上,创建一个新的cv::Mat对象,并将复制的像素数据填充到该对象中。最后,在主机上释放设备内存。 需要注意的是,实现CUDA版本的cv::seamlessClone可能需要一些图像处理和计算机视觉的知识,以及对CUDA编程模型的理解。同时,需要具备使用CUDA编程环境和库函数的能力。 ### 回答3: cv::seamlessClone函数是OpenCV中用于图像无缝融合的函数。要使用CUDA代码实现类似的功能,可以参考以下步骤: 1. 从输入图像和目标图像中读取数据,并将其分配到CUDA设备的全局内存中。 2. 创建一个与输入图像和目标图像大小相同的空白图像作为输出图像,并将其分配到CUDA设备的全局内存中。 3. 在CUDA设备上为输入图像、目标图像和输出图像分配相应的内存空间。 4. 使用CUDA核函数对输入图像和目标图像进行处理,计算图像的梯度(通过Sobel算子或其他方法),并将结果存储在CUDA设备内存中。 5. 使用CUDA核函数对输出图像进行处理,将输入图像和目标图像的梯度信息以及融合参数(比如像素权重)进行计算,并在输出图像中生成无缝融合的效果。 6. 将输出图像从CUDA设备的内存复制到主机内存,以便进一步处理或保存。 7. 释放CUDA设备内存中的图像数据和其他资源。 通过以上步骤,就可以用CUDA代码实现类似于cv::seamlessClone函数的功能,实现图像的无缝融合。但是具体的实现需要根据具体的需求和使用情况来进行一些调整和优化,以提高算法的效率和准确性。
评论 6
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值