CUDA加速向量乘法及结果分析

对两个向量作点乘，可以更改向量维数N，block，thread等参数

#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
#include <time.h>
const int N = 8192;//向量维数

const int threadnum = 32;//使用的线程数

//单block 分散归约 
template <typename T>
__global__ void multi(T *a, T *b, T *c, int n)
{
    __shared__ T tmp[threadnum];
    const int nThreadIdX = threadIdx.x; //线程ID索引号
    const int nBlockDimX = blockDim.x; // 一个block内开启的线程总数
    int nTid = nThreadIdX;
    double dTemp = 0.0;
    while (nTid < n)
    {
        dTemp += a[nTid] * b[nTid];//将结果累加到一个数上
        nTid += nBlockDimX;
    }
    tmp[nThreadIdX] = dTemp; // 将每个线程中的内积放入到共享内存中
    __syncthreads(); // 同步操作，即等所有线程内上面的操作都执行完

    int i=2, j=1;
    while(i <= threadnum)
    {
        if (nThreadIdX%i == 0)
        {
            tmp[nThreadIdX] += tmp[nThreadIdX + j];
        }
        __syncthreads();
        i *= 2;
        j *= 2;
    }
    if (0 == nThreadIdX)
    {
        c[0] = tmp[0];
    }
}


int main()
{
    float a[N], b[N];
    float c = 0;
    for(int i=0; i<N; ++i) // 为数组a、b赋值
    {
        a[i] = i * 1.0;
        b[i] = 1.0;
    }
    float *d_a = NULL, *d_b = NULL, *d_c = NULL;
    cudaMalloc(&d_a, N *sizeof(float));
    cudaMemcpy(d_a, a, N *sizeof(float), cudaMemcpyHostToDevice);

    cudaMalloc(&d_b, N *sizeof(float));
    cudaMemcpy(d_b, b, N *sizeof(float), cudaMemcpyHostToDevice);

    cudaMalloc(&d_c, sizeof(float));
    //计时
    clock_t start,end;
    start = clock();
    
    //计算
    multi<<<1, threadnum>>>(d_a, d_b, d_c, N);
    
    end = clock();
    float time = (float)(end - start) / CLOCKS_PER_SEC;
    printf("总时间为: %f\n",time);
    cudaMemcpy(&c, d_c, sizeof(float), cudaMemcpyDeviceToHost);
        std::cout << c << std::endl;
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    //printf("hello world\n");
    return 0;
}

CmakeLists如下

CMAKE_MINIMUM_REQUIRED(VERSION 2.8)
PROJECT(basis)
FIND_PACKAGE(CUDA REQUIRED)
CUDA_ADD_EXECUTABLE(basis main2.cu)
TARGET_LINK_LIBRARIES(basis)

想查看程序的profile时，发现cuda的nvprof工具用不了，原因是nvprof不支持比较新的显卡设备，并且提示使用Nsight system进行GPU tracing。

网上找了些Nsight system的命令，有些报错了，目前不知道原因，但是用以下命令倒是能显示出类似nvprof的profile结果

nsys profile -t cuda ./basis

basis是make后生成的可执行文件名，该命令结果如下

这个生成的.qdrep文件可以用Nsight system打开（需要在Nvidia官网注册账号后下载），也可以在命令行直接打开.qdrep文件

nsys stats ./report1.qdrep

要是直接复制该指令，执行出错了，就删掉report1.qdrep，从上个命令的结果中拷贝report1.qdrep，再粘贴
结果如下

可以看到核函数的调用时间开销，还有cudamemcpy和cudafree的时间开销，以及memcpy host to device和 device to host的时间对比。