OSlab

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);

运行结果：