对两个向量作点乘,可以更改向量维数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的时间对比。