CUDA学习2-编程部分

最新推荐文章于 2024-07-17 00:57:19 发布

Dodo·D·Caster

最新推荐文章于 2024-07-17 00:57:19 发布

阅读量329

点赞数 1

分类专栏： CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/qq_45347768/article/details/112594723

版权

CUDA 专栏收录该内容

3 篇文章 0 订阅

订阅专栏

CUDA编程

函数声明

host：主机端，通常指CPU

device：设备端，通常指GPU（数据可并行）

kernel：数据并行处理函数，在主机端调用kernel可以在设备端创建大量轻量级线程

host 和 device拥有各自的存储器

CUDA编程包括主机端和设备端两部分代码

	执行位置	调用位置
`__device__float DeviceFunc()`	device	device
`__global__void KernelFunc()`	device	host
`__host__float HostFunc()`	host	host

__global__定义一个kernel函数

入口函数，CPU上调用，GPU上执行
必须返回void

__device__and__host__可以同时使用

global和device函数

尽量少用递归
不要用静态变量
少用malloc
小心通过指针调用

向量数据类型

同时使用于host和device代码

通过函数make_<type name>构造

//构造
int2 i2 = make_int2(1, 2);
float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f);

//访问
int x = i2.x;
int y = i2.y;

数学函数

包含很多常用的数学函数，其中有一个比较特殊：

Intrinsic function 内建函数

仅面向device设备端
更快但精度降低
以__为前缀，如__exp, __log, __sin, __pow, ...

CUDA核函数（kernels）

在N个不同的CUDA线程上并行执行

//定义kernel
_global_ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    //...
    //在N个不同的CUDA线程上并行执行
    VecAdd<<<1, N>>>(A, B, C);
}

线程层次（Thread Hierarchies）

块索引：blockIdx

维度：blockDim

一维、二维或三维

//通过grid和block坐标计算线程id
int threadID = blockIdx.x * blockDim.x + threadIdx.x;
//用线程id从输入读入元素
float x = input[threadID];
//在输入数据上执行函数：数据可并行
float y = func(x);
//用线程id存储输出结果
output[threadID] = y;

//单线程块
//定义kernel
_global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    //...
    //在N*N*1个不同的CUDA线程上并行执行
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    VecAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

//多线程块
//定义kernel
_global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
    int i = threadIdx.x * blockDim.x + threadIdx.x;
    int j = threadIdx.y * blockDim.y + threadIdx.y;
    if(i<N && j<N)
    	C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    //...
    //并行执行
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N/threadsPerBlock.x, N/threadsPerBlock.y);
    VecAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}
/*
N = 32
i = [0,1] * 16 + [0,15]
*/

线程同步

块内线程可以同步

调用__syncthreads创建一个barrier栅栏
每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续指令

Mds[i] = Md[j];
__sncthreads();
func(Mds[i], Mds[i+1]);

同步有时候会导致死锁
if(someFunc())
{
 __syncthreads();
}
else
{
 __syncthreads();
}
比如这里，可能会因为需要同步的进程彼此分开了导致谁也等不到谁从而产生死锁。

线程调度

warp：块内的一组线程（如32个）

线程调度要求warp内所有线程都必须同步。

超出限制后线程数会因为block的减少而减少

内存模型

寄存器（registers）

每个线程专用
快速、片上、可读写

局部存储器（local memory）

存储于global memory
- 每个线程私有
用于存储自动变量数组

共享存储器（shared memory）

每个块
快速、片上、可读写
全速随机访问

全局存储器（global memory）

长延时（100个周期）
片外、可读写
随机访问影响性能
host主机端可读写

常量存储器（constant memory）

短延时、高带宽、当所有线程访问同一位置时只读
存储于global memory但只有缓存
host主机端可读写
容量64kb

变量声明	存储器	作用域	生命期
必须是单独的自动变量而不能是数组	register	thread	kernel
自动变量数组	local	thread	kernel
`__shared__int sharedVar`	shared	block	kernel
`__device__int globalVar`	global	grid	application
`__constant__int constantVar`	constant	grid	application

host可以通过以下函数访问

cudaGetSymbolAddress()
cudaGetSymbolSize()
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()

Constants变量必须在函数外声明

CUDA内存传输

主机端可以从设备端往返传输数据

Global memory 全局存储器

Constant memory 常量存储器

cudaMalloc()：在设备端分配global memory

cudaFree()：释放存储空间

float *Md;	//指向设备端上的一个存储空间
int size = Width * Width * sizeof(float);

cudaMalloc((void**)&Md, size);
//...
cudaFree(Md);

cudaMemcpy()：内存传输

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
//参数：目的地址 源地址 大小 传输方向

host to host
host to device
device to host
device to device

例子：矩阵相乘

//CPU实现
void MatrixMulOnHost(float* M, float* N, float* P, int width)
{
    for(int i=0; i<width; ++i)
        for(int j=0; j<width; ++j)
        {
            float sum = 0;
            for(int k=0; k<width; ++k)
            {
                float a = M[i * width + k];
                float b = N[k * width + j];
                sum += a*b;
            }
            p[i * width + j] = sum;
        }
}

//cuda算法框架（3布）
int main(void)
{
    //1.管理整个内存，为数据分配空间，将数据拷贝到GPU上
    
    //2.在GPU上并行处理计算
    
    //3.将结果拷贝回CPU
}

//GPU实现
void MatrixMulOnDevice(float* M, float* N, float* P, int Width)
{
    int size = Width * Width * sizeof(float);
    //1.管理整个内存，为数据分配空间，将数据拷贝到GPU上
    //分配输入
    cudaMalloc(Md, size);
    cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
    cudaMalloc(Nd, size);
    cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);
    cudaMalloc(Pd, size);
    
    //2.在GPU上并行处理计算
    _global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)
    {
        //访问一个matrix，采用二维block
        int tx = threadIdx.x;
        int ty = threadIdx.y;
        
        //每个kernel线程计算一个输出
        float Pvalue = 0;
        
        //计算
        for(int k=0; k<Width; ++k)
        {
            float Mdelement = Md[ty*Md.width + k];
            float Ndelement = Nd[k*Nd.width + tx];
            Pvalue += Mdelement + Ndelement;
        }
        
        Pd[ty*Width + tx] = Pvalue;
    }
    
    //3.将结果拷贝回CPU
    //1个block含width*width个线程
    dim3 dimBlock(WIDTH, WIDTH);
    dim3 dimGrid(1, 1);
    MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd);
    //传送数据
    cudaMemcpy(Pd, P, size, cudaMemcpyDeviceToHost);
    //释放
    cudaFree(Md);
    cudaFree(Nd);
    cudaFree(Pd);
}