第2.4课 官方库的使用
主要的官方库如上，调用时候可能出现类似“无法解析的外部符号cublasShutdown，该符号在函数main中被引用”这样的错误，需要在vs属性->链接器->输入，里面加入对应得lib库，即cublas.lib,cufft,lib,curand.lib等。
使用库时可查看官方说明https://docs.nvidia.com/cuda/，但是有人说官方的代码有一些错误跑不通，自行留意，满是英文我才看不懂，都是CSDN搜的 ：D
着重介绍curand为矩阵赋初值。有两种方式，
// 需要头文件curand_kernel.h, 利用__device__的随机函数
__global__ void Matrix_init_gpu(float* a, int N) {
int x = threadIdx.x + blockDim.x * blockIdx.x;
curandState state;
long seed = N;
curand_init(seed, x, 0, &state);
if (x < N) a[x] = curand_uniform(&state);
}
//需要头文件curand.h, 利用curand对gpu上的数组进行初始化
void Matrix_init(float* a, int N) {
curandGenerator_t gen;
curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MRG32K3A);
curandSetPseudoRandomGeneratorSeed(gen, 11ULL);
curandGenerateUniform(gen, a, N); // 指定算法生成,生成0-1的随机数
}常用的是下面的方面，用CPU直接调用即可，也就是main函数直接用就行。前三行照着写，不用了解太多，最后一行指定算法生成：
cudaGenerate(curandGenerator_t g,unsigned int * p,size_t N);生成伪随机序列，返回的随机数是32bit的无符号整型。
cudaGenerateLongLong(curandGenerator_t g,unsigned long long * p,size_t N);生成伪随机序列，返回的随机数是64bit的无符号整型。
cudaGenerateUniform(curandGenerator_t g,float * p,size_t N);生成伪随机序列，返回的随机数是0-1的单精度浮点型。
cudaGenerateNormal(curandGenerator_t g,float * p,size_t N,float mean,float stddev);按照指定的均值和方差生成生成伪随机序列。
下面两种方法都用一下，给出代码实例：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "curand.h"
#include "curand_kernel.h"
#include <stdio.h>
#include <iostream>
#include<iomanip>
using namespace std;
// 需要头文件curand_kernel.h, 利用__device__的随机函数
__global__ void Matrix_init_gpu(float* a, int N) {
int x = threadIdx.x + blockDim.x * blockIdx.x;
curandState state;
long seed = N;
curand_init(seed, x, 0, &state);
if (x < N) a[x] = curand_uniform(&state);
}
//需要头文件curand.h, 利用curand对gpu上的数组进行初始化
void Matrix_init(float* a, int N) {
curandGenerator_t gen;
curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MRG32K3A);
curandSetPseudoRandomGeneratorSeed(gen, 11ULL);
curandGenerateUniform(gen, a, N); // 指定算法生成,生成0-1的随机数
}
int main()
{
int m = 16;
int N = m * m;
float* p_da, * p_db, * p_h;
p_h = (float*)malloc(N * sizeof(float)); //在主机侧声明内存空间
cudaMalloc((void**)&p_da, N * sizeof(float));
cudaMalloc((void**)&p_db, N * sizeof(float));
Matrix_init_gpu << <4, 64 >> > (p_da, N);
Matrix_init(p_db, N);
cudaMemcpy(p_h, p_da, N * sizeof(float), cudaMemcpyDeviceToHost); //将随机数传送到主机
printf("\np_da: \n");
for (int i = 0; i < N; i++) { //输出结果
if (i % m == 0)printf("\n");
cout << setw(10) << p_h[i] << " ";
}
cudaMemcpy(p_h, p_db, N * sizeof(float), cudaMemcpyDeviceToHost); //将随机数传送到主机
printf("\np_db: \n");
for (int i = 0; i < N; i++) { //输出结果
if (i % m == 0) printf("\n");
cout << setw(10) << p_h[i] << " ";
}
cudaFree(p_da);
cudaFree(p_db);
free(p_h); //释放主机侧内存空间
}
再说一下，共享内存的事情，上节课说了，共享内存比global memory快几十倍，怎么用共享内存呢？只需要在声明变量的时候前面加上__shared__前缀，当然需要在GPU端函数用。那么用共享内存来计算矩阵乘法的代码就变成了：
// 通过共享内存，进行矩阵的乘法运算
__global__ void Matrix_multi(float* a, float* b, float* c, int m) {
int x = threadIdx.x;
int y = threadIdx.y;
__shared__ float a_shared[256];
__shared__ float b_shared[256];
if (x < m && y < m) {
a_shared[y * m + x] = a[y * m + x];
b_shared[y * m + x] = b[y * m + x];
}
__syncthreads();
float output = 0;
if (x < m && y < m) {
for (int i = 0; i < m; i++) {
output += a_shared[y * m + i] * b_shared[i * m + x];
}
c[y * m + x] = output;
}
}这样的话运行得更快，在数据量大的时候提速很明显。要注意线程的同步，_syncthreads()的运用。
第3课 常见模型
这些模型用cuda解决比较方便，提速明显
1.映射（一对一）
实例：
__global__ void square(float* a, int N) {
int id = threadIdx.x + blockDim.x * blockIdx.x;
if (id < N) {
float s = a[id];//减少读次数
a[id] = s * s;
}
}2.聚合（多对一）
典型：卷积计算。难
为了避免一个block的边界点无法计算卷积的值，有以下两种方法：
1、设置n*n大小的block，把对应坐标的矩阵数据读取到共享内存里，然后其中的(n-2)*(n-2)个线程进行卷积运算
2、设置n*n大小的block，分两次读取矩阵，把(n+2)*(n+2)的矩阵数据读取到内存里，然后进行卷积计算
大家能想明白是什么意思吗？我想明白但是代码不会写，下节课给代码。哎！！！！
3.分散（一对多）
4.转置（一对一）
5.排序
例子：奇偶排序。通过比较数组中相邻的（奇偶）位置数字对，若其排序顺序错误则交换。下一步对所有偶奇位置数字对进行检查并交替。直到结束。
用cuda可以将排序的时间复杂度降至O(N)。
代码：这个函数需要狠狠地理解一下。待排序数组是一个有256个float数的数组。对数组进行排序的时候需要注意线程的同步，接下来开始介绍排序具体操作。
共有256个线程对整个256长度的数组分别进行256次排序，
0号线程调用：进入for循环，第一次j=0,idx=0，对0、1位置上的数字俩排了序，等其他线程排好序。第二次，i=1,j=1,idx=1，对1、2位置上的数字俩排了序，等其他线程排好序。第三次：i=2,j=,idx=0，对0、1位置排了序，等其他线程排好序....如此循环256次。
那么可想而知，1号线程应该是循环对2、3/3、4位置块排序。
...
就这样循环256次，把这个数组排序好了。
这里还用到了动态规划共享内存，前缀extern，后面不用指定数组长度，共享内存比较小，所以最好使用动态规划。extern __shared__ float s_a[];
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "curand.h"
#include "curand_kernel.h"
#include <stdio.h>
#include <iostream>
using namespace std;
// 需要头文件curand_kernel.h, 利用__device__的随机函数
__global__ void Matrix_init_gpu(float* a, int N) {
int x = threadIdx.x + blockDim.x * blockIdx.x;
curandState state;
long seed = N;
curand_init(seed, x, 0, &state);
if (x < N) a[x] = curand_uniform(&state);
}
//需要头文件curand.h, 利用curand对gpu上的数组进行初始化
void Matrix_init(float* a, int N) {
curandGenerator_t gen;
curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MRG32K3A);
curandSetPseudoRandomGeneratorSeed(gen, 11ULL);
curandGenerateUniform(gen, a, N); // 指定算法生成,生成0-1的随机数
}
// 通过共享内存，进行排序
__global__ void sort(float* a, int N) {
int x = threadIdx.x;
extern __shared__ float s_a[];//动态规划
s_a[x] = a[x];
__syncthreads();
for (int i = 0; i < N; i++)
{
int j = i % 2;
int idx = 2 * x + j;
if (idx + 1 < N && s_a[idx] < s_a[idx + 1]) {
float tmp = s_a[idx];
s_a[idx] = s_a[idx + 1];
s_a[idx + 1] = tmp;
}
__syncthreads(); //不要放if里，否则会失效
}
a[x] = s_a[x];
__syncthreads();
}
int main()
{
int m = 10; //输出时每行元素个数
int N = 256;
//curandGenerator_t gen; //生成随机数变量
float* p_d, * p_h, * p_hs;
p_h = (float*)malloc(N * sizeof(float)); //在主机侧声明内存空间
p_hs = (float*)malloc(N * sizeof(float));
cudaMalloc((void**)&p_d, N * sizeof(float)); //GPU侧声明随机数存储缓冲器的内存空间
Matrix_init(p_d, N); //随机数对数组初始化
cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost); //将随机数传送到主机
printf("\np_da: \n");
for (int i = 0; i < N; i++) { //输出结果
if (i % m == 0)printf("\n");
cout << " " << p_h[i] << " ";
}
// 排序算法
sort << <1, N, N * sizeof(float) >> > (p_d, N);
//这里的第三个参数指的是共享内存的大小
cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost); //将随机数传送到主机
printf("\np_d: \n");
for (int i = 0; i < N; i++) { //输出结果
if (i % m == 0) printf("\n");
cout << " " << p_h[i] << " ";
}
cudaFree(p_d); //释放GPU侧内存空间
free(p_h); //释放主机侧内存空间
free(p_hs);
}接下来再做一个训练，做累计求和。
假设这时候有一组数组，把他内部所有值加起来的累加值是多少进行一个计算，那么这个时候可以用0号线程对1、2块累加，1号对3、4,2号对5、6...加完了之后，0号线程再对加完后的1、2块累加........这次累加需要用的线程块是上次的一半，另一半线程不用工作了。
综上，用cuda可以对累加求和实现提速。
代码中，由于共享内存较小，所以运用了两次核函数，先算出一个中间函数，再对这个由中间值组成的中间函数进行累加。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "curand.h"
#include "curand_kernel.h"
#include <stdio.h>
#include <iostream>
#include <time.h>
using namespace std;
// 需要头文件curand_kernel.h, 利用__device__的随机函数
__global__ void Matrix_init_gpu(float *a, int N) {
int x = threadIdx.x + blockDim.x * blockIdx.x;
curandState state;
long seed = N;
curand_init(seed, x, 0, &state);
if (x < N) a[x] = curand_uniform(&state);
}
//需要头文件curand.h, 利用curand对gpu上的数组进行初始化
void Matrix_init(float *a, int N) {
curandGenerator_t gen;
curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MRG32K3A);
curandSetPseudoRandomGeneratorSeed(gen, 11ULL);
curandGenerateUniform(gen, a, N); // 指定算法生成,生成0-1的随机数
}
// 通过共享内存，数组进行求和
__global__ void reduce_kernel(float *d_in, float *d_out) {
int myid = threadIdx.x + blockDim.x * blockIdx.x;
int tid = threadIdx.x;
int blockid = blockIdx.x;
extern __shared__ float sdata[];
sdata[tid] = d_in[myid]; // ** 把d_in[myid]赋值到sdata[tid]
__syncthreads();
for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
if (tid == 0) {
d_out[blockid] = sdata[0];
}
}
void reduce(float *d_in, float *d_mid, float *d_out, int array_size) {
int ThreadsPerBlock = 1024;
int threadnum = ThreadsPerBlock;
int blocks = array_size / ThreadsPerBlock;
reduce_kernel << < blocks, threadnum, threadnum * sizeof(float) >> > (d_in, d_mid);
reduce_kernel << <1, threadnum, threadnum * sizeof(float) >> > (d_mid, d_out);
}
int main()
{
int array_size = 1 << 20;
int array_bytes = array_size * sizeof(float);
//curandGenerator_t gen; //生成随机数变量
float *d_in, *d_mid, *d_out, *h;
float sum = 0;
h = (float *)malloc(array_bytes); //在主机侧声明内存空间
cudaMalloc((void **)&d_in, array_bytes); //GPU侧声明随机数存储缓冲器的内存空间
cudaMalloc((void **)&d_mid, 1024);
cudaMalloc((void **)&d_out, 1);
Matrix_init(d_in, array_size);
cudaMemcpy(h, d_in, array_bytes, cudaMemcpyDeviceToHost); //将随机数传送到主机
clock_t t_start_cpu = clock();
for (int i = 0; i < array_size; i++)
{
sum += h[i];
}
cout << "sum_cpu:" << sum << endl;
cout << "CPU耗时：" << (double)(clock() - t_start_cpu) / CLOCKS_PER_SEC << "毫秒" << endl;
clock_t t_start_gpu = clock();
reduce(d_in, d_mid, d_out, array_size);
cudaDeviceSynchronize();
clock_t cpy_time = clock();
cudaMemcpy(h, d_out, array_bytes, cudaMemcpyDeviceToHost); //将随机数传送到主机
cout << "sum_gpu:" << h[0] << endl;
cout << "GPU耗时：" << (double)(clock() - t_start_gpu) / CLOCKS_PER_SEC << "毫秒" << endl;
cout << "数组传输的耗时" << (double)(clock() - cpy_time) / CLOCKS_PER_SEC << "毫秒" << endl;
free(h);
cudaFree(d_in);
cudaFree(d_mid);
cudaFree(d_out);
return 0;
}最后CPU和GPU计算的结果会不一样，不要问我为啥，我不知道...有人说是小数位保留的问题，可我算出来差的也太大了，我也不知道为啥。真的不懂。
2039
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
