OSlab
B站链接:
https://www.bilibili.com/video/BV1kU4y1m7QW?share_source=copy_web
Linux环境实践作业
在Windows Terminal的Ubuntu 20.04版命令行界面下,提前写好了两个hello.c和makefile文件,通过make指令编译hello.c文件,并在命令行中输出hello,world!greetings from <姓名>-<学号>,再通过make clean指令删除make.i make.S make.o make等文件。
Intel OneAPI实践作业
oneAPI实践练习
通过ssh工具,连接devcloud:
ssh devcloud
克隆oneAPI_course文件夹:
git clone https://github.com/pengzhao-intel/oneAPI_course
进入code文件夹:
cd oneAPI_course/code/
问题1:改写gemm_basic代码26,27行,利用work group和local work item的坐标来计算
修改前gemm_basic.cpp源码26~27行::
int row = index.get_global_id(0);
int col = index.get_global_id(1);
编译并运行gemm_basic.cpp:
dpcpp gemm_basic.cpp -o gemm_basic
vim run1
qsub -lnodes=1:ppn=2:gpu -d . run1
run1内容:
./gemm_basic
qstat查看状态:
运行结果:
修改gemm_basic.cpp文件:
vim gemm_basic.cpp
修改内容:
int row = index.get_local_id(0) + index.get_group(0)*block_size;
int col = index.get_local_id(1) + index.get_group(1)*block_size;
运行结果:
问题二:修改程序输入数据的大小,设定为M=N=K=2000,修改程序,并使其通过正确性测试
修改内容:
int errCode = gemm(2000, 2000, 2000, 4, 10, my_gpu_queue);
运行结果:
结果:GPU和CPU计算时间大大增加
问题三:测试不同tile_X和tile_Y大小下矩阵计算的性能,并分析原因
修改内容:
#define tileY n
#define tileX n
n取2:
n取4:
n取8:
n取16:
n | GPU | CPU |
---|---|---|
2 | 8.585604 | 61.146653 |
4 | 4.609323 | 49.536103 |
8 | 4.581816 | 49.518492 |
16 | 29.404795 | 51.445425 |
原因分析:
tile:对计算部分的划分,数字越大,划分越细
tile小,任务太多,会影响性能;tile大,性能过剩。猜测:tile和block_size相等时性能最好。
结果分析:tile在4-8之间,计算机性能最佳。
oneAPI编程作业: 计算CrossEntropy
代码实现
引入数学头文件,满足Softmax函数对数、指数的计算:
#include <math>
GPU内核函数:
首先降维处理,计算时降为一维;同时,将X看做为[K][M]的二维矩阵,sumX也为[K][M]的二维矩阵。使用block给矩阵分块。
auto local_ndrange = range<2>(block_size, block_size);
auto global_ndrange = range<2>(grid_rows, grid_cols);
parallel_for并行处理函数体
h.parallel_for<class k_name_tt>(sycl::nd_range<2>(global_ndrange, local_ndrange), [=](sycl::nd_item<2> index) {
int row = index.get_global_id(0);
int col = index.get_global_id(1);
C[row*N+col] = - B[row*M*N+mask[row*N+col]*N+col]*weight[row*N+col];
});
最后实现GPU的多线程运算。
考虑到最优性能,将三维矩阵、二维矩阵降维到一维。
例如,最后一步中的cC[K][N]和weight[K][N]的下标变换为i*N+k
for(int k=0;k<N;k++)
{
cC[i*N+k]=-cB[i*M*N+mask[i*N+k]*N+k]*weight[i*N+k];
}
CPU是单线程的运算,相较与GPU实现更为简单。
Verify函数:
验证cpu和gpu的运算结果,结果误差小于千分之一就正确,否则报错
int verify(float *cpu_res, float *gpu_res, int length){
int err = 0;
for(int i = 0; i < length; i++) {
if( fabs(cpu_res[i] - gpu_res[i]) > 1e-3) {
err++;
printf("\n%lf, %lf", cpu_res[i], gpu_res[i]);
}
}
return(err);
}
Cross函数:
分配地址空间
原问题中的X转换,分配内存空间大小为K* M* N,用一维数组来表示三维
auto X = malloc_shared<float>(K*M*N, q);
auto Y = malloc_shared<float>(K*M*N, q);
用一维数组表示二维,loss和loss_host保存GPU和CPU的运算结果。
auto loss = malloc_shared<float>(K*N, q); //GPU
auto loss_host = malloc_host<float>(K*N, q);//CPU
auto mask = malloc_shared<int>(K*N,q);
auto weight = malloc_shared<float>(K*N,q);