h5做招聘网站,wordpress固定,上海建设牌电动三轮官方网站,南宁网页制作步骤一 基本函数 在GPU上开辟空间#xff0c;无论定义的数据是float还是int ,还是****gpu_int,分配空间的函数都是下面固定的形式 (void**) 1.函数定义#xff0c;global void 是配套使用的#xff0c;是在GPU上定义#xff0c;也就是GPU上执行#xff0c;CPU上调用的函数…一 基本函数 在GPU上开辟空间无论定义的数据是float还是int ,还是****gpu_int,分配空间的函数都是下面固定的形式 (void**) 1.函数定义global void 是配套使用的是在GPU上定义也就是GPU上执行CPU上调用的函数因为CPU不能识别GPU上运算得到的结果也就是说在CPU上调用这个函数是不可能存在return的结果的所以没有返回值 #include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h// __是GOU函数的标识符加了__标识这个函数是在GPU上被调用的//这个函数在GPU上是无法调用的所谓在GPU上就是定义的在GPU上的函数内调用因为该函数当前是在CPU上
int add_one(int a)
{return a 1;
}//这个函数是被定义在GPU上由GPU本身调用的函数所以加了这个__devide__标识之后就可以在__global__ void show(int *a)函数里面调用该函数
__device__ int add_one(int a)
{return a 1;
}//但是加了__devide__标识之后由于表示只能在GPU上调用该函数所以在mian函数里面如果想要实现a[i] add_one(a[i]);是不可以的
//所以需要再加一个标识符
__host__ __device__ int add_one(int a)
{return a 1;
}__global__ void show(int *a)
{for (int i 0; i 10; i){a[i] add_one(a[i]);printf( %d , a[i]);}printf(\n);
}#include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h// __是GOU函数的标识符加了__标识这个函数是在GPU上被调用的//这个函数在GPU上是无法调用的所谓在GPU上就是定义的在GPU上的函数内调用因为该函数当前是在CPU上
//int add_one(int a)
//{
// return a 1;
//}
//
这个函数是被定义在GPU上由GPU本身调用的函数所以加了这个__devide__标识之后就可以在__global__ void show(int *a)函数里面调用该函数
//__device__ int add_one(int a)
//{
// return a 1;
//}//但是加了__devide__标识之后由于表示只能在GPU上调用该函数所以在mian函数里面如果想要实现a[i] add_one(a[i]);是不可以的
//所以需要再加一个标识符
__host__ __device__ int add_one(int a)
{return a 1;
}__global__ void show(int *a)
{for (int i 0; i 10; i){// a[i] add_one(a[i]);printf( %d , a[i]);}printf(\n);
}__global__ void int_gpu(int* a)
{for (int i 0; i 10; i){a[i] 100;}
}int main()
{int cpu[10] { 10, 10, 10, 10, 10, 10, 10, 10, 10, 10};//在GPU上分配空间存储CPU上的数据int* gpu_int;cudaMalloc((void**)gpu_int, 10*sizeof(int)); //将指针指向GPU的一个内存地址show 1, 1 (gpu_int);//一个网格里面只有一个block,也就是只有一个线程//GPU上数组初始化cudaMemset(gpu_int, 0, 10 * sizeof(int));// 将CPU上的数据拷贝到GPU上cudaMemcpy(gpu_int, cpu, 10*sizeof(int), cudaMemcpyHostToDevice); //将数据从GPU拷贝到CPU,同时指定拷贝的长度// 在CPU上调用GPU上的函数进行计算定义在GPU上的函数就是在GPU上进行运算show 1, 1 (gpu_int);//一个网格里面只有一个block,也就是只有一个线程int_gpu 1, 1 (gpu_int);show 1, 1 (gpu_int);// 将GPU上的数据拷贝到cPU上cudaMemcpy(cpu, gpu_int, 10*sizeof(int), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度cudaFree(gpu_int);cudaDeviceSynchronize();printf( cpu \n);for (int i 0; i 10; i){printf( %d , cpu[i]);}return 0;
}二 gird, block, thread 之间的关系和理解 
1 同步的使用时机 2 尽量避免直接从globla memory上频繁读写可以将数据拷贝到share memory再进行对同一数据频繁的读写操作
local memory数据读取是很快的 global memeory,从这里读取数据是很慢的 share memory, 是block自己共享的每个block自己可以读自己内部的数据一个block内部的线程可以访问自己block的share memory
3
threadIdx是指当前线程在当前线程块里面是排几号是相对block来说的 blockIdx是当前线程块在当前grid里面x这个维度是第几个线程块 blockDim是维度的意思 下面这种情况对应的就是grid是2维block是二维的情况含有Dim的是对整个grid维度或者blcok维度的定义也就是有几个 矩阵运算
案例1 矩阵维度为20*20获取矩阵位置坐标x,y并返回xy #include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostreamusing namespace std;//20*20的矩阵 坐标x,y位置赋值xy__device__ int coord_int(int x, int y)
{return x y;
}__global__ void Matrix_init(int *a, int m, int n)
{int x blockIdx.x * blockDim.x threadIdx.x;int y blockIdx.y * blockDim.y threadIdx.y;if (x m y n){a[y * n x] coord_int(x, y);}
}void show(int* a, int m, int n)
{for (int i 0; i n; i){for ( int j 0; j m; j){cout a[i * m j] ;}cout endl;}
}int main()
{int* gpu_int;cudaMalloc((void**)gpu_int, 400 * sizeof(int));int cpu_int[400] { 0 };show(cpu_int, 20, 20);dim3 blockdim(8, 8); //线程数要大于400dim3 griddim(3, 3);//为什么数组要开辟一个新的空间且要用指针而数据不用因为数据编译的时候是可以直接读的不需要开辟空间Matrix_init griddim, blockdim (gpu_int, 20, 20);cudaMemcpy(cpu_int, gpu_int, 400 * sizeof(int), cudaMemcpyDeviceToHost);cudaFree(gpu_int);show(cpu_int, 20, 20);return 0;
}案例2 矩阵乘法和矩阵加法
#include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostreamusing namespace std;//矩阵20*20
__host__ __device__ int add_one(int a)
{return a 1;
}__global__ void Matrix_add(int* a, int* b, int* c, int m, int n)
{int x blockIdx.x * blockDim.x threadIdx.x;int y blockIdx.y * blockDim.y threadIdx.y;if (x m y n){c[y * n x] a[y * n x] b[y * n x];}
}//矩阵乘法前一个矩阵的行*后一个矩阵的列然后相加
//只适用与方阵
__global__ void Matrix_multi(int* a, int* b, int* c, int m)
{int x blockIdx.x * blockDim.x threadIdx.x;int y blockIdx.y * blockDim.y threadIdx.y;int output 0;//c[y * m x] 0; if (x m y m){for (int i 0; i m; i){output a[y * m i] b[i * m x];//c[y * m x] a[y * m i] b[i * m x]; 这里相当于全局内存反复读取c数组会导致速度慢很多}c[y * m x] output;}
}//cpu上的函数
void show(int* a, int m, int n)
{for (int i 0; i m; i){for (int j 0; j n; j){// a[i] add_one(a[i]);cout a[i * m j] ;}cout endl;}}__global__ void int_gpu(int* a,int m, int n)
{for (int i 0; i m; i){for (int j 0; j n; j){a[i * m j] 10;}}
}//__global__ void my_gpu_multi(int* a, int* b, int* c, int m)
//{
// __share__ int
//}int main()
{int cpu[400] { 0 };show(cpu, 20, 20);//在GPU上分配空间存储CPU上的数据int* gpu_int;cudaMalloc((void**)gpu_int, 400 * sizeof(int)); //将指针指向GPU的一个内存地址int* gpu_add;cudaMalloc((void**)gpu_add, 400 * sizeof(int)); //将指针指向GPU的一个内存地址int* gpu_multi;cudaMalloc((void**)gpu_multi, 400 * sizeof(int)); //将指针指向GPU的一个内存地址//GPU上数组初始化cudaMemset(gpu_int, 0, 400 * sizeof(int));cudaMemset(gpu_add, 0, 400 * sizeof(int));cudaMemset(gpu_multi, 0, 400 * sizeof(int));// 将CPU上的数据拷贝到GPU上cudaMemcpy(gpu_int, cpu, 400 * sizeof(int), cudaMemcpyHostToDevice); //将数据从GPU拷贝到CPU,同时指定拷贝的长度cudaMemcpy(gpu_add, cpu, 400 * sizeof(int), cudaMemcpyHostToDevice); //将数据从GPU拷贝到CPU,同时指定拷贝的长度cudaMemcpy(gpu_multi, cpu, 400 * sizeof(int), cudaMemcpyHostToDevice); //将数据从GPU拷贝到CPU,同时指定拷贝的长度// 在CPU上调用GPU上的函数进行计算定义在GPU上的函数就是在GPU上进行运算dim3 blockdim(8, 8);dim3 griddim(3, 3);int_gpu griddim, blockdim (gpu_int, 20, 20);Matrix_add griddim, blockdim (gpu_int, gpu_int, gpu_add, 20,20);Matrix_multi griddim, blockdim (gpu_int, gpu_int, gpu_multi, 20);// 将GPU上的数据拷贝到cPU上cudaMemcpy(cpu, gpu_int, 400 * sizeof(int), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度show(cpu, 20, 20);//一个网格里面只有一个block,也就是只有一个线程cudaMemcpy(cpu, gpu_add, 400 * sizeof(int), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度show(cpu, 20, 20);//一个网格里面只有一个block,也就是只有一个线程cudaMemcpy(cpu, gpu_multi, 400 * sizeof(int), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度show(cpu, 20, 20);//一个网格里面只有一个block,也就是只有一个线程cudaFree(gpu_int);cudaFree(gpu_add);cudaFree(gpu_multi);cudaDeviceSynchronize();return 0;
}案例3 用共享内存实现矩阵乘法 #include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream
#include curand.h
#include curand_kernel.husing namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16//GPU上初始化,N是
__global__ void gpu_initial(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);
}//在CPU上使用随机数初始化N是要初始化的数组长度
void cpu_initial(float *a, int N)
{curandGenerator_t gen;curandCreateGenerator(gen, CURAND_RNG_PSEUDO_DEFAULT); //告诉电脑要用什么样的方法生成随机数curandSetPseudoRandomGeneratorSeed(gen, 11ULL); //指定随机数种子curandGenerateUniform(gen, a, N); //curandGenerateUniform 生成均匀分布的0-1的随机数让a里面的数据从1-N的数据都进行赋值
}void show(float* a, int m, int n)
{for (int i 0; i m; i){for (int j 0; j n; j){// a[i] add_one(a[i]);cout a[i * m j] ;}cout endl;}}//矩阵乘法前一个矩阵的行*后一个矩阵的列然后相加
//只适用与方阵
__global__ void Matrix_multi(float* a, float* b, float* c, int m)
{int x threadIdx.x;int y threadIdx.y;int output 0;if (x m y m){for (int i 0; i m; i){output a[y * m i] * b[i * m x];}c[y * m x] output;}
}__global__ void my_gpu_multi(float* a, float* b, float* c, int m)
{int x threadIdx.x;int y threadIdx.y;__shared__ float a_share[256];__shared__ float b_share[256];if (x m y m){a_share[y * m x] a[y * m x];b_share[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_share[y*m i] * b_share[i*m x];}c[y * m x] output;}}fighting
共享内存大小不足
共享内存的大小由 __shared__ float a_share[256]; 和 __shared__ float b_share[256]; 决定。如果 m 值大于 16那么共享内存大小可能不足因为 16x16 256。如果 m 大于 16应该调整共享内存大小。
同步问题
在 my_gpu_multi 中使用了 __syncthreads() 来确保所有线程都完成了数据拷贝。但如果有些线程的计算还没完成就进入下一步可能会导致不一致的结果。
线程边界检查
如果 m 的值较大而 block 尺寸 (m, m) 超出了 GPU 硬件的限制可能会导致一些线程未能正确启动导致结果不一致。int main()
{int m 6;int N m * m;float* p_d, * p_da, * p_db, * p_h;//cpu上开辟空间p_h (float*)malloc(N * sizeof(float));//GPU上开辟空间cudaMalloc((void**)p_d, N * sizeof(float)); //将指针指向GPU的一个内存地址cudaMalloc((void**)p_da, N * sizeof(float)); //将指针指向GPU的一个内存地址cudaMalloc((void**)p_db, N * sizeof(float)); //将指针指向GPU的一个内存地址//数组初始化使用两种不同的方式进行初始化gpu_initial 16, 16 (p_da, N); //直接在GPU上初始化cpu_initial(p_db, N); //在CPU上初始化,但是由于是调用的Gpu上的函数进行的初始化所以最后还是等价于在GPU上初始化的cudaMemcpy(p_h, p_da, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度printf(\n p_da);for (int i 0; i N; i){if (i % m 0) printf(\n );cout p_h[i] ;}cudaMemcpy(p_h, p_db, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度printf(\n p_db);for (int i 0; i N; i){if (i % m 0) printf(\n );cout p_h[i] ;}printf(\n );//share memory只能在一个block上使用所以只能定义一个blockdim3 blockdim(m, m);//Matrix_multi 1, blockdim (p_da, p_db, p_d, m); 将GPU上的数据拷贝到cPU上//cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度//printf(funtuon Matrix_multi: ------------\n );//show(p_h, m, m);//一个网格里面只有一个block,也就是只有一个线程my_gpu_multi 1, blockdim (p_da, p_db, p_d, m);cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度printf(funtuon my_gpu_multi: ------------\n );show(p_h, m, m);cudaFree(p_da);cudaFree(p_db);cudaFree(p_d);free(p_h);return 0;
}
常用官方库的使用
1. cuda案例
cuda函数说明官网 //Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include stdio.h
#include stdlib.h
#include math.h
#include cuda_runtime.h
#include cublas_v2.h
#define M 6
#define N 5
#define IDX2F(i,j,ld) ((((j)-1)*(ld))((i)-1))static __inline__ void modify(cublasHandle_t handle, float* m, int ldm, int n, int p, int q, float alpha, float beta) {cublasSscal(handle, n - q 1, alpha, m[IDX2F(p, q, ldm)], ldm);cublasSscal(handle, ldm - p 1, beta, m[IDX2F(p, q, ldm)], 1);
}int main(void) {// cudaError_t可以判断cuda有没有错误如果定义的cudaStat不是cudaSucess就会提示错误信息位置cudaError_t cudaStat;cublasStatus_t stat;cublasHandle_t handle;int i, j;float* devPtrA;float* a 0;a (float*)malloc(M * N * sizeof(*a));if (!a) {printf(host memory allocation failed);return EXIT_FAILURE;}for (j 1; j N; j) {for (i 1; i M; i) {a[IDX2F(i, j, M)] (float)((i - 1) * N j);}}cudaStat cudaMalloc((void**)devPtrA, M * N * sizeof(*a));if (cudaStat ! cudaSuccess) {printf(device memory allocation failed);free(a);return EXIT_FAILURE;}stat cublasCreate(handle);if (stat ! CUBLAS_STATUS_SUCCESS) {printf(CUBLAS initialization failed\n);free(a);cudaFree(devPtrA);return EXIT_FAILURE;}stat cublasSetMatrix(M, N, sizeof(*a), a, M, devPtrA, M);if (stat ! CUBLAS_STATUS_SUCCESS) {printf(data download failed);free(a);cudaFree(devPtrA);cublasDestroy(handle);return EXIT_FAILURE;}modify(handle, devPtrA, M, N, 2, 3, 16.0f, 12.0f);stat cublasGetMatrix(M, N, sizeof(*a), devPtrA, M, a, M);if (stat ! CUBLAS_STATUS_SUCCESS) {printf(data upload failed);free(a);cudaFree(devPtrA);cublasDestroy(handle);return EXIT_FAILURE;}cudaFree(devPtrA);cublasDestroy(handle);for (j 1; j N; j) {for (i 1; i M; i) {printf(%7.0f, a[IDX2F(i, j, M)]);}printf(\n);}free(a);return EXIT_SUCCESS;
}2.cuda自带函数实现矩阵乘法运算 cublasSgemm lad是A的行数ldb是B的行数 //Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include stdio.h
#include stdlib.h
#include math.h
#include cuda_runtime.h
#include cublas_v2.h
#include ctime
#include iostreamusing namespace std;int main()
{srand(time(0));// C A*Bint M 2; //矩阵A的行矩阵C的行int N 3; //矩阵A的列矩阵B的行int K 4; //矩阵B的列矩阵C的列float* h_A (float*)malloc(sizeof(float) * M * N);float* h_B (float*)malloc(sizeof(float) * N * K);float* h_C (float*)malloc(sizeof(float) * M * K);//******************************* 矩阵初始化cout A ----------- endl;for (int i 0; i M * N; i){h_A[i] rand() % 10;cout h_A[i] ;if ((i 1) % N 0)cout endl;}cout endl;cout B ----------- endl;for (int i 0; i N * K; i){h_B[i] rand() % 10;cout h_B[i] ;if ((i 1) % K 0)cout endl;}cout endl;//******************************* 在GPU上开辟空间存储数据float *d_A, *d_B, *d_C;cudaMalloc( (void**)d_A, sizeof(float) * M * N);cudaMalloc((void**)d_B, sizeof(float) * N * K);cudaMalloc((void**)d_C, sizeof(float) * M * K);//******************************* 将数据从CPU拷贝到GPUcudaMemcpy(d_A, h_A, sizeof(float) * M * N, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, sizeof(float) * N * K, cudaMemcpyHostToDevice);float alpha 1;float beta 0;cublasHandle_t handle;cublasCreate(handle);cublasSgemm(handle,CUBLAS_OP_N, //数据不转置CUBLAS_OP_N,K, //矩阵B的列M, //矩阵A的行N, //矩阵A的列alpha,d_B,K,d_A,N,beta,d_C,K);cudaMemcpy(h_C, d_C, sizeof(float) * M * K, cudaMemcpyDeviceToHost);cout c ----------- endl;for (int i 0; i M * K; i){cout h_C[i] ;if ((i 1) % K 0)cout endl;}cout endl;cudaFree(d_A);cudaFree(d_B);cudaFree(d_C);return 0;
}3.cuda自带函数实现矩阵每个数的翻倍 cublasSccal
把间隔设置为1就可以实现把当前矩阵每一个元素都乘以这么一个常量 //Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include stdio.h
#include stdlib.h
#include math.h
#include cuda_runtime.h
#include cublas_v2.h
#include ctime
#include iostreamusing namespace std;int main()
{srand(time(0));// C A*Bint M 2; //矩阵A的行矩阵C的行int N 3; //矩阵A的列矩阵B的行float* h_A (float*)malloc(sizeof(float) * M * N);//******************************* 矩阵初始化cout A ----------- endl;for (int i 0; i M * N; i){h_A[i] rand() % 10;cout h_A[i] ;if ((i 1) % N 0)cout endl;}cout endl;//******************************* 在GPU上开辟空间存储数据float *d_A, *d_B, *d_C;cudaMalloc( (void**)d_A, sizeof(float) * M * N);//******************************* 将数据从CPU拷贝到GPUcudaMemcpy(d_A, h_A, sizeof(float) * M * N, cudaMemcpyHostToDevice);float alpha 2.2;cublasHandle_t handle;cublasStatus_t stat;//函数用于初始化CUBLAS库并创建一个CUBLAS上下文。上下文被表示为一个cublasHandle_t类型的句柄。//handle是一个指向cublasHandle_t类型的指针cublasCreate函数会将创建的句柄存储在这个位置。//如果初始化成功stat将返回CUBLAS_STATUS_SUCCESS否则会返回一个错误码。stat cublasCreate(handle); if (stat ! CUBLAS_STATUS_SUCCESS) {printf(CUBLAS initialization failed\n);return EXIT_FAILURE;}stat cublasSscal(handle,6, alpha, d_A,1);if (stat ! CUBLAS_STATUS_SUCCESS) {printf(CUBLAS scaling failed\n);return EXIT_FAILURE;}cudaMemcpy(h_A, d_A, sizeof(float) * M * N, cudaMemcpyDeviceToHost);cout c ----------- endl;for (int i 0; i M * N; i){cout h_A[i] ;if ((i 1) % N 0)cout endl;}cout endl;cudaFree(d_A);return 0;
}4. 傅里叶变换 https://docs.nvidia.com/cuda/cufft/index.html
cufftExecC2C() and cufftExecZ2Z()是一样的前者是浮点型后者是double后者精度更高 //Example 1. Application Using C and cuBLAS: 1-based indexing
//-----------------------------------------------------------
#include stdio.h
#include stdlib.h
#include math.h
#include cuda_runtime.h
#include cublas_v2.h
#include ctime
#include iostream
#include cufft.husing namespace std;int main()
{const int Nt 256;const int BACTH 1;//BACTH 用户霹雳那个处理一批一维的数据假设数据是512当BATCH 2,则将0-255,256-512作为两个一维信号做FFT变换//定义cufftDoubleComplex类型的指针device_in记录输入值device_out计记录计算结果cufftDoubleComplex *host_in, * host_out, * device_in, * device_out;//这种方式开辟空空间再向GPU传输数据是要快于传统的malloc方式的两者是等价的/* float* host_in (float*)malloc(sizeof(float) * Nt);float* host_out (float*)malloc(sizeof(float) * Nt);*/cudaMallocHost((void**)host_in, sizeof(cufftDoubleComplex) * Nt);cudaMallocHost((void**)host_out, sizeof(cufftDoubleComplex) * Nt);for (int i 0; i Nt; i){host_in[i].x i 1;host_in[i].y i 1;}// ************* 在GPU上开辟空间cudaMalloc((void**)device_in, sizeof(cufftDoubleComplex) * Nt);cudaMalloc((void**)device_out, sizeof(cufftDoubleComplex) * Nt);//从CPU上数据拷贝数据到GPU上cudaMemcpy(device_in, host_in, Nt* sizeof(cufftDoubleComplex), cudaMemcpyHostToDevice);//创建cufft句柄cufftHandle cufftForwrdHandle;cufftPlan1d(cufftForwrdHandle, Nt, CUFFT_Z2Z, BACTH); //传参//CUFFT_Z2Z因为前面定义的数据类型是Double-Complex,所以这里指定的数据类型是这个// typedef enum cufftType_t {// CUFFT_R2C 0x2a, // Real to Complex (interleaved)// CUFFT_C2R 0x2c, // Complex (interleaved) to Real// CUFFT_C2C 0x29, // Complex to Complex, interleaved// CUFFT_D2Z 0x6a, // Double to Double-Complex// CUFFT_Z2D 0x6c, // Double-Complex to Double// CUFFT_Z2Z 0x69 // Double-Complex to Double-Complex//} cufftType;//执行fft正变换cufftExecZ2Z(cufftForwrdHandle, device_in, device_out, CUFFT_FORWARD); //正变换是CUFFT_FORWARD反变换是CUFFT_INVERSE// 从GPU 上数据拷贝数据到CPU上cudaMemcpy(host_out, device_out, Nt * sizeof(cufftDoubleComplex), cudaMemcpyDeviceToHost);//设置输出精度--输出正变换的结果cout 正变换的结果: endl;for (int i 0; i Nt; i){cout host_out[i].x j* host_out[i].y endl;}cudaFree(device_in);cudaFree(device_out);return 0;
}三 常见的可以采用并行的模式
1, 一对一 例如输入x,函数输出y x*x #include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream
#include curand.h
#include curand_kernel.husing namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16//GPU上初始化,N是
__global__ void gpu_initial(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);
}__global__ void square(float *d_out,int N)
{int myID threadIdx.x blockDim.x * blockIdx.x;if (myID N){float data d_out[myID];d_out[myID] data * data;}
}int main()
{int m 6;int N m * m;float* d_out, *h_in, *h_out;//cpu上开辟空间h_in (float*)malloc(N * sizeof(float));h_out (float*)malloc(N * sizeof(float));//GPU上开辟空间cudaMalloc((void**)d_out, N * sizeof(float)); //将指针指向GPU的一个内存地址//数组初始化使用两种不同的方式进行初始化//直接在GPU上初始化, 4 是 numBlocks表示启动了 4 个块block,16 是 numThreadsPerBlock表示每个块中有 16 个线程thread)//这种初始化方式适用于一维数组gpu_initial 4, 16 (d_out, N); cudaMemcpy(h_in, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);printf(\n h_in);for (int i 0; i N; i){if (i % m 0) printf(\n );cout h_in[i] ;}square 4, 16 (d_out, N);//将数据从GPU拷贝到CPU,同时指定拷贝的长度cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);printf(\n h_out);for (int i 0; i N; i){if (i % m 0) printf(\n );cout h_out[i] ;}cudaFree(d_out);free(h_in);free(h_out);return 0;
}
2. 卷积操作 我写的
#include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream
#include curand.h
#include curand_kernel.husing namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16//GPU上初始化,N是
__global__ void gpu_initial(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);
}void show(float* a, int m, int n)
{for (int i 0; i m; i){for (int j 0; j n; j){// a[i] add_one(a[i]);cout a[i * m j] ;}cout endl;}}//设置n*n大小的block把对应坐标的矩阵数据读取到共享内存里然后其中的(n-2)*(n-2)个线程进行卷积运算
__global__ void Conv_method1(float* p_d, float*kernel, int block_size, int kernel_Len)
{int x threadIdx.x; //现在是一个一维的一个blockint y threadIdx.y;extern __shared__ float s_pd[];if (x block_size y block_size){s_pd[y * block_size x] p_d[y * block_size x];//printf(s_pd[y * block_size x] %f \n , s_pd[y * block_size x]);}__syncthreads();//for (int i 0; i kernel_Len; i) {// printf(kernel[%d] %f, i, kernel[i]);//}float out 0;if (x 0 y 0 x block_size-1 y block_size-1){out s_pd[(y-1)* block_size x - 1] * kernel[0] s_pd[(y - 1) * block_size x] * kernel[1] s_pd[(y - 1) * block_size x1] * kernel[2] s_pd[y * block_size x - 1] * kernel[3] s_pd[y * block_size x] * kernel[4] s_pd[y * block_size x 1] * kernel[5] s_pd[(y1)* block_size x - 1] * kernel[6] s_pd[(y 1) * block_size x] * kernel[7] s_pd[(y 1) * block_size x 1] * kernel[8];//printf( x %d,y %d,out %f \n,x,y,out);}p_d[y * block_size x] out;__syncthreads();
}int main()
{int m 6;int N m*m;int kernelLen 9;float* p_d, * h_in, * h_out;float *h_kenel, * d_kernel;//cpu上开辟空间h_in (float*)malloc(N * sizeof(float));h_out (float*)malloc(N * sizeof(float));h_kenel (float*)malloc(9 * sizeof(float));//GPU上开辟空间cudaMalloc((void**)p_d, N * sizeof(float)); //将指针指向GPU的一个内存地址cudaMalloc((void**)d_kernel, 9 * sizeof(float)); //将指针指向GPU的一个内存地址 使用循环赋值 printf(h_kenel \n );float values[] { 1, 1, 1, 1, -8, 1, 1,1,1};for (int i 0; i kernelLen; i) {h_kenel[i] values[i];cout h_kenel[i] ;}printf( \n );cudaMemcpy(d_kernel, h_kenel, kernelLen * sizeof(float), cudaMemcpyHostToDevice);gpu_initial 16, 16 (p_d, N); //直接在GPU上初始化cudaMemcpy(h_in, p_d, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度printf(h_in \n );show(h_in, m, m);dim3 griddim(1, 1);dim3 blockdim(m, m);Conv_method1 griddim, blockdim, sizeof(float)* N (p_d, d_kernel, m, kernelLen);//将数据从GPU拷贝到CPU,同时指定拷贝的长度cudaMemcpy(h_out, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);printf(\n h_out \n );show(h_out,m,m);cudaFree(p_d);free(h_in);free(h_out);return 0;
}别人写的
#include cuda_runtime.h
#include device_launch_parameters.h#include iostream
#include fstream
#include iomanip
#include time.husing namespace std;// 调用CUDA函数并检查是否出现错误
#define CUDA_CALL(x) {\const cudaError_t a (x);\if (a ! cudaSuccess) { \fprintf(stderr, \nCUDA Error: %s (err_num%d)\nfile %s, line %d\n, cudaGetErrorString(a), a, __FILE__, __LINE__);\cudaDeviceReset(); exit(1);\} \
}// 调用核函数并检查是否出现错误
__host__ void cuda_error_check(const char* kernelName) {if (cudaPeekAtLastError() ! cudaSuccess) {printf(\n%s %s\n, kernelName, cudaGetErrorString(cudaGetLastError()));cudaDeviceReset();exit(1);}
}// 设置参数
const int TILE_W 32; // block的x维大小
const int TILE_H 32; // block的y维大小
const int DATA_W 320; // 输入矩阵的x维大小
const int DATA_H 640; // 输入矩阵的y维大小
const int KERNEL_RADIUS 1; // 卷积核的半径// 卷积核写入GPU常量内存中
__constant__ int KERNEL[2 * KERNEL_RADIUS 1][2 * KERNEL_RADIUS 1]
{ 1, 1, 1,
1, -8, 1,
1, 1, 1 };// 通过共享内存进行卷积运算
__global__ void convolution(float* dst, float* src) {// 线程在所在block中的x,y坐标int tidx threadIdx.x;int tidy threadIdx.y;// 线程应该读取到shared memory中对应的矩阵元素坐标int readx (blockDim.x - 2 * KERNEL_RADIUS) * blockIdx.x (threadIdx.x - KERNEL_RADIUS);int ready (blockDim.y - 2 * KERNEL_RADIUS) * blockIdx.y (threadIdx.y - KERNEL_RADIUS);// 除去不需要加载内存的地方if (readx DATA_W KERNEL_RADIUS || ready DATA_H KERNEL_RADIUS)return;// 声明block的共享内存__shared__ float src_s[TILE_H][TILE_W];// 把当前block需要处理的区域读取到shared memory输入矩阵周围的记为0if (readx 0 readx DATA_W ready 0 ready DATA_H) {src_s[tidy][tidx] src[ready * DATA_W readx];}else {src_s[tidy][tidx] 0;}// 同步block中所有线程保证共享内存完全读入矩阵__syncthreads();// 卷积计算float output 0;int kernel_w 2 * KERNEL_RADIUS 1;if (tidx blockDim.x - 2 * KERNEL_RADIUS readx DATA_W - KERNEL_RADIUS tidy blockDim.y - 2 * KERNEL_RADIUS ready DATA_H - KERNEL_RADIUS) {for (int i 0; i kernel_w; i) {for (int j 0; j kernel_w; j) {output src_s[tidy j][tidx i] * KERNEL[j][i];}}// 写入dst对应坐标dst[(ready KERNEL_RADIUS) * DATA_W (readx KERNEL_RADIUS)] output;}
}int main() {const int INPUTSIZE DATA_H * DATA_W;printf(---------- initilizing ----------\n);clock_t tt clock();// CPU输入输出矩阵的声明float* h_src (float*)malloc(INPUTSIZE * sizeof(float));float* h_dst (float*)malloc(INPUTSIZE * sizeof(float));// 输入矩阵中元素全部设为1for (int i 0; i DATA_W; i) {for (int j 0; j DATA_H; j) {h_src[i j * DATA_W] (float)1;}}// 将输入矩阵输出ofstream ofs(input_output.txt);for (int j 0; j DATA_H; j) {for (int i 0; i DATA_W; i) {ofs setw(5) h_src[i j * DATA_W];}ofs \n;}ofs \n;// 设定使用第一个GPUCUDA_CALL(cudaSetDevice(0));// GPU输入输出矩阵的声明float* d_src 0;float* d_dst 0;// 初始化CUDA_CALL(cudaMalloc(d_src, INPUTSIZE * sizeof(float)));CUDA_CALL(cudaMalloc(d_dst, INPUTSIZE * sizeof(float)));CUDA_CALL(cudaMemcpy(d_src, h_src, INPUTSIZE * sizeof(float), cudaMemcpyHostToDevice));CUDA_CALL(cudaMemcpy(d_dst, h_dst, INPUTSIZE * sizeof(float), cudaMemcpyHostToDevice));// 输出内存初始化所需时间printf(Initializaion time(ms) : %f\n, ((float)clock() - tt) / CLOCKS_PER_SEC);printf(---------- calculating ----------\n);// 设定核函数中grid和block的参数//dim3 gridDim (11, 22);/*gridDim定义了网格的维度。在这个例子中网格有两个维度x 维度有 11 个块y 维度有 22 个块*/dim3 gridDim { (DATA_W (TILE_W - 2 * KERNEL_RADIUS - 1)) / (TILE_W - 2 * KERNEL_RADIUS),(DATA_H (TILE_H - 2 * KERNEL_RADIUS - 1)) / (TILE_H - 2 * KERNEL_RADIUS) };printf(---------- calculating ----(DATA_W (TILE_W - 2 * KERNEL_RADIUS - 1)) / (TILE_W - 2 * KERNEL_RADIUS) %d, (DATA_H (TILE_H - 2 * KERNEL_RADIUS - 1)) / (TILE_H - 2 * KERNEL_RADIUS) %d--\n, (DATA_W (TILE_W - 2 * KERNEL_RADIUS - 1)) / (TILE_W - 2 * KERNEL_RADIUS),(DATA_H (TILE_H - 2 * KERNEL_RADIUS - 1)) / (TILE_H - 2 * KERNEL_RADIUS));dim3 blockDim { TILE_W, TILE_H };// 调用核函数并计时cudaEvent_t start, stop;cudaEventCreate(start);cudaEventCreate(stop);cudaEventRecord(start, 0);convolution gridDim, blockDim (d_dst, d_src);// 检查核函数调用是否出现错误cuda_error_check(convolution_shared_memory);// CPU等待GPU完成核函数的计算CUDA_CALL(cudaDeviceSynchronize());// 输出核函数调用时长cudaEventRecord(stop, 0);cudaEventSynchronize(stop);float elapsedTime;cudaEventElapsedTime(elapsedTime, start, stop);printf(Kernel time(ms) : %f\n, elapsedTime);cudaEventDestroy(start);cudaEventDestroy(stop);// CPU获取GPU核函数计算结果CUDA_CALL(cudaMemcpy(h_dst, d_dst, INPUTSIZE * sizeof(float), cudaMemcpyDeviceToHost));// 输出卷积后的结构for (int j 0; j DATA_H; j) {for (int i 0; i DATA_W; i) {ofs setw(5) h_dst[i j * DATA_W];}ofs \n;}ofs.close();// 释放GPU内存cudaFree(d_src);cudaFree(d_dst);//cudaFree(KERNEL);// 释放CPU内存free(h_src);free(h_dst);// 清空重置GPUCUDA_CALL(cudaDeviceReset());printf(Total calculation time(ms) : %f\n, ((float)clock() - tt) / CLOCKS_PER_SEC);}3. 4. 5 动态共享内存的初始化和使用 #include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream
#include curand.h
#include curand_kernel.husing namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16//GPU上初始化,N是
__global__ void gpu_initial(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);
}//使用共享内存
__global__ void sortFun(float* p_d, int N)
{int x threadIdx.x; //现在是一个一维的一个block//定义动态内存,动态内存的长度是在初始化的时候 sortFun 1, N, sizeof(float)*N (d_in, N);sizeof(float)*N 指定共享内存的字节大小extern __shared__ float temp_d[]; temp_d[x] p_d[x];__syncthreads(); //相当于是把所有数据都加载到temp_d之后才开始后面的排序操作//奇偶排序长度为N的数组需要排序N次所以这里的for循环里面的N是对排序次数的循环而不是数据for (int i 0; i N; i) {int j i % 2; //先奇偶比然后下一次才是偶奇比也就是偶数次是奇偶比奇数次是偶奇比int idx 2 * x j;if (idx 1 N temp_d[idx] temp_d[idx 1]){float tep temp_d[idx];temp_d[idx] temp_d[idx 1];temp_d[idx 1] tep;}__syncthreads(); //每排序一次要等所有数据都判断完毕才进行下一轮的排序}p_d[x] temp_d[x];__syncthreads(); //所有数据都排序完毕且结束程序结束
}int main()
{int m 6;int N m;float* p_d, *h_in, *h_out;//cpu上开辟空间h_in (float*)malloc(N * sizeof(float));h_out (float*)malloc(N * sizeof(float));//GPU上开辟空间cudaMalloc((void**)p_d, N * sizeof(float)); //将指针指向GPU的一个内存地址 使用循环赋值 //int values[] { 1, 5, 4, 2, 6, 3};//for (int i 0; i N; i) {// h_in[i] values[i];//}//cudaMemcpy(p_d, h_in, N * sizeof(float), cudaMemcpyHostToDevice);gpu_initial 16, 16 (p_d, N); //直接在GPU上初始化cudaMemcpy(h_in, p_d, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度printf(h_in \n );for (int i 0; i N; i){cout h_in[i] ;}sortFun 1, N, sizeof(float)*N (p_d, N);// sort 1, N, N * sizeof(float) (p_d, N);//将数据从GPU拷贝到CPU,同时指定拷贝的长度cudaMemcpy(h_out, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);printf(\n h_out \n );for (int i 0; i N; i){cout h_out[i] ;}cudaFree(p_d);free(h_in);free(h_out);return 0;
}6 !!! 数组求和 #include cuda_runtime.h
#include device_launch_parameters.h#include stdio.h
#include iostream
#include curand.h
#include curand_kernel.husing namespace std;
#define ARRAY_SIZE 16
#define ARRAY_LENGTH 256
#define BLOCK_SIZE 16//GPU上初始化,N是
__global__ void gpu_initial(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);
}void show(float* a, int m, int n)
{for (int i 0; i m; i){for (int j 0; j n; j){// a[i] add_one(a[i]);cout a[i * m j] ;}cout endl;}
}__global__ void reduce_kernel (float *data_in, float *data_out)
{//现在是N个1维的blockint threadIDinAllBlock threadIdx.x blockDim.x*blockIdx.x; int threadIDinOneBlock threadIdx.x;int blockID blockIdx.x;extern __shared__ float s_d[];//共享内存是在一个blcok,但是传入的数据是分散到1024个block上s_d[threadIDinOneBlock] data_in[threadIDinAllBlock];__syncthreads();//for (int i 0; i N; i)//{// printf(here s_d[%d] %f \n,i, s_d[i]);//}//printf( \n);//对半求和前一半数据的第一位数和后半数据的第一位数相加。。。//这里是算完一个blockfor (int s blockDim.x/2; s 0; s 1) //s 1 是除以2的意思,数据依次是1024,512,256.。。。。{if (threadIDinOneBlock s){s_d[threadIDinOneBlock] s_d[threadIDinOneBlock s];}__syncthreads();}//输出1024个blcok的各自的结果if (threadIDinOneBlock 0){data_out[blockID] s_d[0];/*printf(here s_d[0] %f \n, s_d[0]);*/printf(here data_out[%d] %f \n, blockID, data_out[blockID]);}
}void cpu_reduce(float *d_in, float* d_mid, float* d_out, int N)
{int threadNum N;int blocks N / threadNum;//printf(blocks %d \n, blocks);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 m 50;int N m*m;float* d_in,*d_mid, *d_out, * h_in, * h_out;//cpu上开辟空间h_in (float*)malloc(N * sizeof(float));h_out (float*)malloc(N * sizeof(float));//GPU上开辟空间cudaMalloc((void**)d_in, N * sizeof(float)); //将指针指向GPU的一个内存地址cudaMalloc((void**)d_mid, m * sizeof(float));//一个block的大小cudaMalloc((void**)d_out, N * sizeof(float));gpu_initial m, m (d_in, N); //直接在GPU上初始化cudaMemcpy(h_in, d_in, N * sizeof(float), cudaMemcpyDeviceToHost); //将数据从GPU拷贝到CPU,同时指定拷贝的长度clock_t t_start_cpu clock();float sum 0;printf(in mian \n);for (int i 0; i N; i){sum h_in[i];//cout h_in[i] ;}cout \n;cout sum_ cpu sum endl;cout time in cpu (double)(clock() - t_start_cpu) / CLOCKS_PER_SEC endl;cout \n;cout ----------------------- \n;clock_t t_start_gpu clock();cpu_reduce(d_in, d_mid,d_out,N);cudaDeviceSynchronize();cudaMemcpy(h_out, d_out,N * sizeof(float), cudaMemcpyDeviceToHost);cout SUM gpu : h_out h_out[0] endl;cout time in GPU (double)(clock() - t_start_gpu) / CLOCKS_PER_SEC endl;cudaFree(d_in);cudaFree(d_out);cudaFree(d_mid);free(h_in);free(h_out);return 0;
}
