- 虽然是作业,但是我也觉得了解一下cuda还是很有必要的,之前本科想做硬件加速时也不知道应该先学学cuda是怎么工作的。
- Linux上安装cuda后, cuda其实提供了很多例程, 如下图. 但是还是感觉太复杂了,所以这里总结几个简单的cuda程序.
- 需要注意的是使用了cuda的c程序后缀名应为
.cu
, 同时编译器为nvcc
,如nvcc check_device.cu -o check_deviec
1 查看设备
在cuda安装目录下的./include/cuda_runtime_api.h
中可以看到对结构体cudaDeviceProp
的说明.
// check_device.cu
#include <stdio.h>
//#include <cuda_runtime.h>
int main(int argc, char **argv)
{
int deviceCount;
cudaGetDeviceCount(&deviceCount);
printf("Device count: %d\n", deviceCount);
int device;
for(device = 0; device < deviceCount; ++device)
{
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp,device);
printf("Device id: %d, \n\tDevice name: %s\n", device, deviceProp.name);
printf("\tCompute capability %d.%d.\n",deviceProp.major,deviceProp.minor);
printf("\tNum of Stream Multiprocess(SM): %d\n", deviceProp.multiProcessorCount);
printf("\t每个线程块的最大线程数: %d\n", deviceProp.maxThreadsPerBlock);
printf("\t每个SM的最大线程数: %d, 最大线程束数: %d\n",
deviceProp.maxThreadsPerMultiProcessor, deviceProp.maxThreadsPerMultiProcessor/32);
}
return 0;
}
2 向量加法
- cuda编写简单的程序,不着重考虑性能的话只需要了解一下grid和block的概念,以及
__global__
和__device__
关键字,参考1中写的很简明了.编写核函数时思维和在cpu上编程非常不同, 要考虑其如何并行, **即每个thread不能依赖于其他thread的结果。**但实际代码其实只需稍作改动就可以。
// vector_adder.cu
// 程序有点问题,运行后显存不释放
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
// 每个thread只做一个元素的加法
__global__ void add_1( int *a, int *b, int *c, int num)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
c[idx] = a[idx] + b[idx];
// Fermi以后的架构,可以在核函数中调用printf
// printf("%d\n",c[idx]);
}
// thread数量有限,当数据量大的时候, 就需要每个thread做多个元素的加法.
// 这里的代码是类比参考2 CUDA编程入门极简教程,这个代码逻辑没问题,也能并行计算,但我认为不符合cuda的编程思想。应该按照参考1 线程块调度的思想来写程序。也就是核函数中只关注一次计算就可以了。
__global__ void add_2( int *a, int *b, int *c, int num)
{
// 获取thread在整个grid中的索引.
int idx = threadIdx.x + blockIdx.x * blockDim.x;
// 计算thread总数
int stride = blockDim.x * gridDim.x;
// 每个thread处理 num/stride 个元素
for (int i = idx; i < num; i += stride)
{
c[i] = a[i] + b[i];
}
}
/* 我们最好是参考这种形式来写,然后分配足够多的block和thread,显然之后大于实际可用的thread数量,但是gpu会按照block进行调度。当gpu中的计算单元完成核函数的计算后就会被释放,此时等候的block就得以分配到SM上。
*/
/*
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
*/
int main(int argc, char **argv)
{
// 向量长度和字节数
int N = 20;
int bytes = N*sizeof(int);
int pass = 1;
int *dev_c1;
int *dev_c2;
int *dev_b;
int *dev_a;
// 申请host内存,即CPU接的内存
int *vec_a = (int *)malloc(bytes);
int *vec_b = (int *)malloc(bytes);
int *vec_c1 = (int *)malloc(bytes);
int *vec_c2 = (int *)<