并行矩阵乘法
描述
编写⼀个基于oneAPI的C++/SYCL程序来执行矩阵乘法操作。需要考虑大尺寸矩阵的乘法操作以及不同线程之间的数据依赖关系。通常在实现矩阵乘法时，可以使用块矩阵乘法以及共享内存来提高计算效率。
分析
利用基于SYCL的编程模型在GPU上实现矩阵乘法的计算，步骤如下：
-
分配内存：在主机端分配内存空间用于存储输⼊矩阵和输出矩阵，同时在GPU端分配内存空间用于存储相应 的输入和输出数据。
-
数据传输：将输入矩阵数据从主机端内存传输到GPU端内存中。
-
核函数调用：在SYCL中，矩阵乘法的计算通常会在GPU上使用核函数来实现并行计算。核函数 会分配线程块和线程来处理不同的数据块。
-
并行计算：在核函数中，每个线程负责计算输出矩阵的⼀个单独的元素。为了最大限度地利用 GPU的并行计算能力，通常会使用⼆维线程块和线程网格的方式来处理矩阵的乘法计算。
-
数据传输：计算完成后，将输出矩阵数据从GPU端内存传输回主机端内存中，以便进⼀步处理或 分析。 在并行计算矩阵乘法时，可以利用线程块和线程的层次结构来优化计算。通过合理划分矩阵数据并利用共享内存来减少全局内存访问的次数，可以大幅提高计算效率。此外，还可以利用GPU上的多个计算单元并执行行矩阵乘法，进⼀步提高计算速度。
代码
#include <chrono>
#include <iostream>
#include <CL/sycl.hpp>
#define random_float() (rand() / double(RAND_MAX))
using namespace std;
using namespace sycl;
#define tileY 2
#define tileX 2
double gpu_kernel(float *A, float *B, float *C,
int M, int N, int K,
int BLOCK, sycl::queue &q) {
auto grid_rows = M / tileY;
auto grid_cols = N / tileX;
auto local_ndrange = range<2>(BLOCK, BLOCK);
auto global_ndrange = range<2>(grid_rows, grid_cols);
double duration = 0.0f;
auto e = q.submit([&](sycl::handler &h) {
h.parallel_for<class k_name_t>(
sycl::nd_range<2>(global_ndrange, local_ndrange), [=](sycl::nd_item<2> index) {
int row = tileY * index.get_global_id(0);
int col = tileX * index.get_global_id(1);
float sum[tileY][tileX] = {0.0f};
float subA[tileY] = {0.0f};
float subB[tileX] = {0.0f};
// core computation
for (int k = 0; k < N; k++) {
// read data to register
for(int m = 0; m < tileY; m++) {
subA[m] = A[(row + m) * N + k];
}
for(int p = 0; p < tileX; p++) {
subB[p] = B[k * N + p + col];
}
for (int m = 0; m < tileY; m++) {
for (int p = 0; p < tileX; p++) {
sum[m][p] += subA[m] * subB[p];
}
}
} //end of K
for (int m = 0; m < tileY; m++) {
for (int p = 0; p < tileX; p++) {
C[(row + m) * N + col + p] = sum[m][p];
}
}
});
});
e.wait();
duration += (e.get_profiling_info<info::event_profiling::command_end>() -
e.get_profiling_info<info::event_profiling::command_start>()) /1000.0f/1000.0f;
return(duration);
}
double cpu_kernel(float *cA, float *cB, float *cC, int M, int N, int K) {
double duration = 0.0;
std::chrono::high_resolution_clock::time_point s, e;
// Single Thread Computation in CPU
s = std::chrono::high_resolution_clock::now();
for(int i = 0; i < M; i++) {
for(int j = 0; j < N; j++) {
float sum = 0.0f;
for(int k = 0; k < K; k++) {
sum += cA[i * K + k] * cB[k * N + j];
}
cC[i * N + j] = sum;
}
}
e = std::chrono::high_resolution_clock::now();
duration = std::chrono::duration<float, std::milli>(e - s).count();
return(duration);
}
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);
}
int gemm(const int M,
const int N,
const int K,
const int block_size,
const int iterations,
sycl::queue &q) {
cout << "Problem size: c(" << M << "," << N << ") ="
<< " a(" << M << "," << K << ") *"
<< " b(" << K << "," << N << ")\n";
auto A = malloc_shared<float>(M * K, q);
auto B = malloc_shared<float>(K * N, q);
auto C = malloc_shared<float>(M * N, q);
auto C_host = malloc_host<float>(M * N, q);
for(int i=0; i < M * K; i++) {
A[i] = random_float();
}
for(int i=0; i < K * N; i++) {
B[i] = random_float();
}
for(int i=0; i < M * N; i++) {
C[i] = 0.0f;
C_host[i] = 0.0f;
}
double flopsPerMatrixMul
= 2.0 * static_cast<double>(M) * static_cast<double>(N) * static_cast<double>(K);
double duration_gpu = 0.0f;
double duration_cpu = 0.0f;
int warmup = 10;
for (int run = 0; run < iterations + warmup; run++) {
float duration = gpu_kernel(A, B, C, M, N, K, block_size, q);
if(run >= warmup) duration_gpu += duration;
}
duration_gpu = duration_gpu / iterations;
warmup = 2;
for(int run = 0; run < iterations/2 + warmup; run++) {
float duration = cpu_kernel(A, B, C_host, M, N, K);
if(run >= warmup) duration_cpu += duration;
}
duration_cpu = duration_cpu / iterations/2;
int errCode = 0;
errCode = verify(C_host, C, M*N);
if(errCode > 0) printf("\nThere are %d errors\n", errCode);
printf("\nGEMM size M = %d, N = %d, K = %d", M, N, K);
printf("\nWork-Group size = %d * %d, tile_X = %d, tile_Y = %d", block_size, block_size, tileX, tileY);
printf("\nPerformance Flops = %lf, \n"
"GPU Computation Time = %lf (ms); \n"
"CPU Computaiton Time = %lf (ms); \n",
flopsPerMatrixMul, duration_gpu, duration_cpu);
free(A, q);
free(B, q);
free(C, q);
free(C_host, q);
return(errCode);
}
int main() {
auto propList = cl::sycl::property_list {cl::sycl::property::queue::enable_profiling()};
queue my_gpu_queue( cl::sycl::gpu_selector{} , propList);
int errCode = gemm(512, 512, 512,
4,
10,
my_gpu_queue);
return(errCode);
}
运行结果
并行排序算法
描述
使用基于oneAPI的C++/SYCL实现⼀个高效的并行归并排序。需要考虑数据的分割和合并以及线程之间的协作。
分析
归并排序是⼀种分治算法，其基本原理是将待排序的数组分成两部分，分别对这两部分进行排序，然后将已排序的子数组合并为⼀个有序数组。可考虑利用了异构并行计算的特点，将排序和合并操作分配给多个线程同时执行，以提高排序效率。具体实现过程如下：将待排序的数组分割成多个较小的子数组，并将这些⼦数组分配给不同的线程块进行处理。每个线程块内部的线程协作完成子数组的局部排序。通过多次迭代，不断合并相邻的有序⼦数组，直到整个数组有序。 在实际实现中，归并排序可使用共享内存来加速排序过程。具体来说，可以利用共享内存来存储临时数据，减少对全局内存的访问次数，从而提高排序的效率。另外，在合并操作中，需要考虑同步机制来保证多个线程之间的数据⼀致性。 需要注意的是，在实际应用中，要考虑到数组大小、线程块大小、数据访问模式等因素，来设计合适的算法和参数设置，以充分利用目标计算硬件GPU的并行计算能力，提高排序的效率和性能。
代码
# include <iostream>
using namespace std;
void swap(double* p, double* q); //函数声明, 交换两个变量的值
void quickSort(double* array, int low, int high); //函数声明, 快速排序
int main(void)
{
double array[] = { 76, 7, 32, 4, 43, 9, 1, 56, 8 };
int len = sizeof(array) / sizeof(array[0]);
quickSort(array, 0, len-1);
cout << "快速排序：";
for (int i = 0; i < len; i++)
cout << array[i] << " ";
cout << endl;
return 0;
}
void swap(double* p, double* q)
{
double temp = *p;
*p = *q;
*q = temp;
return;
}
void quickSort(double* array, int low, int high)
{
if (low >= high) // 排序结束，返回
return;
int i = low; // 指向第一个元素
int j = high; // 指向最后一个元素
double key = array[low]; // 选择一个基准值
while (low < high) // 一次循环表示一轮快排
{
while (low < high && key <= array[high])
--high; //向前寻找
if (key > array[high])
{
swap(&array[low], &array[high]);
++low;
}
while (low < high && key >= array[low])
++low; //向后寻找
if (key < array[low])
{
swap(&array[low], &array[high]);
--high;
}
}
quickSort(array, i, low - 1);
quickSort(array, low + 1, j);
}
运行结果
792
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
