GPU编程笔记2------CUDA

cuda线程块组织结构

gpu 线程组织 分为线程块 和线程 两个层级 其中线程块（block） 对线程进行组织，grid对线程块进行组织。线程块和线程都是三维数据结构。
每个线程块中数量不可超过1024个（三维总数）。
同一个线程块中的线程可以共享内存。

cudaError_t cudaMallocManaged(void ** ,size_t size); //统一内存申请 通过内存迁移方式 使内存即支持gpu又支持cpu

第一个cuda程序 使用1个block 计算add函数

#include<iostream>
#include<math.h>
__global__
void add(int n,float *x,float *y){
        int index = threadIdx.x; //当前线程编号
        int stride = blockDim.x; //总共线程数量 步长

        for(int i=index;i<n;i+= stride) //多线程 for循环 每个线程处理不同的i 
                y[i] = x[i] + y[i];
}


int main(){
        int N = 1<<20;

        float *x,*y;
        cudaMallocManaged(&x,N*sizeof(float));	//统一内存申请
        cudaMallocManaged(&y,N*sizeof(float));	//统一内存申请

        for(int i=0;i<N;i++){
                x[i] = 1.0f;
                y[i] = 2.0f;
        }

        add<<<1,256>>>(N,x,y);		//使用gpu计算 1个线程块中256个线程 并行计算

        cudaDeviceSynchronize();	//等待内核执行完成
        float maxError = -1.0;
        for(int i=0;i<N;i++){
                maxError = fmax(maxError,fabs(y[i]-3.0f));
        }

        std::cout<<maxError<<std::endl;

        cudaFree(x);
        cudaFree(y);

        return 0;
}

修改程序 使用10个block 计算add函数

#include<iostream>
#include<math.h>
__global__
void add(int n,float *x,float *y){
        int index = blockIdx.x * blockDim.x + threadIdx.x; //多个线程块 计算当前线程编号
        int stride = gridDim.x * blockDim.x;		//总共线程数量 步长 gridDim.x 表示使用的线程块数量

        for(int i=index;i<n;i+= stride) //多线程 for循环 每个线程处理不同的i 
                y[i] = x[i] + y[i];
}


int main(){
        int N = 1<<20;

        float *x,*y;
        cudaMallocManaged(&x,N*sizeof(float));	//统一内存申请
        cudaMallocManaged(&y,N*sizeof(float));	//统一内存申请

        for(int i=0;i<N;i++){
                x[i] = 1.0f;
                y[i] = 2.0f;
        }

        add<<<10,256>>>(N,x,y);		//使用gpu计算 1个线程块中256个线程 并行计算

        cudaDeviceSynchronize();	//等待内核执行完成
        float maxError = -1.0;
        for(int i=0;i<N;i++){
                maxError = fmax(maxError,fabs(y[i]-3.0f));
        }

        std::cout<<maxError<<std::endl;

        cudaFree(x);
        cudaFree(y);

        return 0;
}

add 计算时间明显变小了，继续增加线程块数量，发现计算时间不会在提高太多了！！
继续优化 将初始化操作移动到 gpu中进行
添加初始化函数

__global__ 
void init(int n,float *x,float *y){
        int index = blockIdx.x * blockDim.x + threadIdx.x;
        int stride = gridDim.x * blockDim.x; 

        for(int i=index;i<n;i+= stride){
                x[i] = 1.0f;
                y[i] = 2.0f;
        }
}

总耗时又减少了不少，add 仅用了11.424 us （下图）
分析：cudaMallocManaged 申请内存是在gpu上的 如果在cpu上对数据进行操作，cuda会进行内存页的迁移，将gpu内存迁移到cpu。然后调用add函数在gpu上计算，cuda又会将cpu内存迁移到gpu。nvprof 将内存页移动时间计算在add运算时长内。所以耗时较高。