matrix fixed
This commit is contained in:
325
perflab/poly/poly.cu
Normal file
325
perflab/poly/poly.cu
Normal file
@@ -0,0 +1,325 @@
|
||||
/**************************************************************************
|
||||
多项式计算函数。按下面的要求编辑此文件：
|
||||
1. 将你的学号、姓名，以注释的方式写到下面；
|
||||
2. 实现不同版本的多项式计算函数；
|
||||
3. 编辑peval_fun_rec peval_fun_tab数组，将你的最好的答案
|
||||
（最小CPE、最小C10）作为数组的前两项
|
||||
***************************************************************************/
|
||||
|
||||
/*
|
||||
学号：201209054233
|
||||
姓名：夜半加班狂
|
||||
*/
|
||||
|
||||
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <cuda_runtime.h>
|
||||
typedef int (*peval_fun)(int*, int, int);
|
||||
|
||||
typedef struct {
|
||||
peval_fun f;
|
||||
char *descr;
|
||||
} peval_fun_rec, *peval_fun_ptr;
|
||||
|
||||
|
||||
/**************************************************************************
|
||||
Edit this comment to indicate your name and Andrew ID
|
||||
#ifdef ASSIGN
|
||||
Submission by Harry Q. Bovik, bovik@andrew.cmu.edu
|
||||
#else
|
||||
Instructor's version.
|
||||
Created by Randal E. Bryant, Randy.Bryant@cs.cmu.edu, 10/07/02
|
||||
#endif
|
||||
***************************************************************************/
|
||||
|
||||
/*
|
||||
实现一个指定的常系数多项式计算
|
||||
第一次，请直接运行程序，以便获知你需要实现的常系数是啥
|
||||
*/
|
||||
int const_poly_eval(int *not_use, int not_use2, int x)
|
||||
{
|
||||
int result = 0;
|
||||
/* int i;
|
||||
int xpwr = 1; // x的幂次
|
||||
int a[4] = {21,90,42,88};
|
||||
for (i = 0; i <= 3; i++) {
|
||||
result += a[i]*xpwr;
|
||||
xpwr *= x;
|
||||
}
|
||||
*/
|
||||
// 90 = 64 + 32 - 4 - 2
|
||||
// 42 = 32 + 8 + 2
|
||||
// 88 = 64 + 16 + 8
|
||||
int x64,x32,x16,x8,x4,x2;
|
||||
|
||||
x64 = x << 6;
|
||||
x32 = x << 5;
|
||||
x16 = x << 4;
|
||||
x8 = x << 3;
|
||||
x4 = x << 2;
|
||||
x2 = x << 1;
|
||||
result = 21 + x64+x32-x4-x2 + ((x32+x8+x2) + (x64+x16+x8)*x)*x;
|
||||
return result;
|
||||
}
|
||||
|
||||
|
||||
|
||||
/* 多项式计算函数。注意：这个只是一个参考实现，你需要实现自己的版本 */
|
||||
|
||||
/*
|
||||
友情提示：lcc支持ATT格式的嵌入式汇编，例如
|
||||
|
||||
_asm("movl %eax,%ebx");
|
||||
_asm("pushl %edx");
|
||||
|
||||
可以在lcc中project->configuration->Compiler->Code Generation->Generate .asm，
|
||||
将其选中后，可以在lcc目录下面生成对应程序的汇编代码实现。通过查看汇编文件，
|
||||
你可以了解编译器是如何实现你的代码的。有些实现可能非常低效。
|
||||
你可以在适当的地方加入嵌入式汇编，来大幅度提高计算性能。
|
||||
*/
|
||||
|
||||
int poly_eval(int *a, int degree, int x)
|
||||
{
|
||||
int result = 0;
|
||||
int i;
|
||||
int xpwr = 1; /* x的幂次 */
|
||||
// printf("阶=%d\n",degree);
|
||||
for (i = 0; i <= degree; i++) {
|
||||
result += a[i]*xpwr;
|
||||
xpwr *= x;
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
/* CUDA优化的多项式计算函数 - 低CPE版本 */
|
||||
int cuda_poly_eval_low_cpe(int *a, int degree, int x)
|
||||
{
|
||||
// 对于低CPE版本，我们使用CUDA并行计算多项式的各个项
|
||||
// 然后将结果传回主机进行求和
|
||||
|
||||
// 分配设备内存
|
||||
int *d_a, *d_results;
|
||||
cudaError_t err;
|
||||
|
||||
// 分配内存
|
||||
err = cudaMalloc(&d_a, (degree + 1) * sizeof(int));
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
return 0;
|
||||
}
|
||||
|
||||
err = cudaMalloc(&d_results, (degree + 1) * sizeof(int));
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 将系数从主机复制到设备
|
||||
err = cudaMemcpy(d_a, a, (degree + 1) * sizeof(int), cudaMemcpyHostToDevice);
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_results);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 定义CUDA核函数
|
||||
dim3 blockDim(256);
|
||||
dim3 gridDim((degree + 1 + blockDim.x - 1) / blockDim.x);
|
||||
|
||||
// 启动核函数
|
||||
cudaPolyEvalLowCPE<<<gridDim, blockDim>>>(d_a, degree, x, d_results);
|
||||
|
||||
// 检查核函数执行错误
|
||||
err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_results);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 分配主机内存用于结果
|
||||
int *h_results = (int *)malloc((degree + 1) * sizeof(int));
|
||||
if (h_results == NULL) {
|
||||
printf("Memory allocation error\n");
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_results);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 将结果从设备复制回主机
|
||||
err = cudaMemcpy(h_results, d_results, (degree + 1) * sizeof(int), cudaMemcpyDeviceToHost);
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
free(h_results);
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_results);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 在主机上求和
|
||||
int result = 0;
|
||||
for (int i = 0; i <= degree; i++) {
|
||||
result += h_results[i];
|
||||
}
|
||||
|
||||
// 释放内存
|
||||
free(h_results);
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_results);
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
/* CUDA优化的多项式计算函数 - 10阶优化版本 */
|
||||
int cuda_poly_eval_degree10(int *a, int degree, int x)
|
||||
{
|
||||
// 对于10阶多项式，我们可以使用更优化的方法
|
||||
// 使用CUDA并行计算，但针对10阶多项式进行特殊优化
|
||||
|
||||
// 分配设备内存
|
||||
int *d_a, *d_result;
|
||||
cudaError_t err;
|
||||
|
||||
// 分配内存
|
||||
err = cudaMalloc(&d_a, (degree + 1) * sizeof(int));
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
return 0;
|
||||
}
|
||||
|
||||
err = cudaMalloc(&d_result, sizeof(int));
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 将系数从主机复制到设备
|
||||
err = cudaMemcpy(d_a, a, (degree + 1) * sizeof(int), cudaMemcpyHostToDevice);
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_result);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 定义CUDA核函数
|
||||
dim3 blockDim(256);
|
||||
dim3 gridDim(1); // 只需要一个块，因为我们只需要一个结果
|
||||
|
||||
// 启动核函数
|
||||
cudaPolyEvalDegree10<<<gridDim, blockDim>>>(d_a, degree, x, d_result);
|
||||
|
||||
// 检查核函数执行错误
|
||||
err = cudaGetLastError();
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_result);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 获取结果
|
||||
int result;
|
||||
err = cudaMemcpy(&result, d_result, sizeof(int), cudaMemcpyDeviceToHost);
|
||||
if (err != cudaSuccess) {
|
||||
printf("CUDA Error: %s\n", cudaGetErrorString(err));
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_result);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// 释放内存
|
||||
cudaFree(d_a);
|
||||
cudaFree(d_result);
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
/* CUDA核函数 - 低CPE版本 */
|
||||
__global__ void cudaPolyEvalLowCPE(int *a, int degree, int x, int *results)
|
||||
{
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx <= degree) {
|
||||
// 计算x的幂
|
||||
int xpwr = 1;
|
||||
for (int i = 0; i < idx; i++) {
|
||||
xpwr *= x;
|
||||
}
|
||||
|
||||
// 计算这一项的结果
|
||||
results[idx] = a[idx] * xpwr;
|
||||
}
|
||||
}
|
||||
|
||||
/* CUDA核函数 - 10阶优化版本 */
|
||||
__global__ void cudaPolyEvalDegree10(int *a, int degree, int x, int *result)
|
||||
{
|
||||
// 使用共享内存来存储中间结果
|
||||
__shared__ int shared_result;
|
||||
|
||||
// 只有第一个线程初始化共享结果
|
||||
if (threadIdx.x == 0) {
|
||||
shared_result = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
// 每个线程计算一部分项
|
||||
int local_result = 0;
|
||||
int xpwr = 1;
|
||||
|
||||
// 计算x的幂
|
||||
for (int i = 0; i < threadIdx.x; i++) {
|
||||
xpwr *= x;
|
||||
}
|
||||
|
||||
// 计算这一项的结果
|
||||
if (threadIdx.x <= degree) {
|
||||
local_result = a[threadIdx.x] * xpwr;
|
||||
}
|
||||
|
||||
// 使用原子操作累加结果
|
||||
atomicAdd(&shared_result, local_result);
|
||||
|
||||
// 同步所有线程
|
||||
__syncthreads();
|
||||
|
||||
// 只有第一个线程将结果写回全局内存
|
||||
if (threadIdx.x == 0) {
|
||||
*result = shared_result;
|
||||
}
|
||||
}
|
||||
|
||||
/*
|
||||
这个表格包含多个数组元素，每一组元素（函数名字, "描述字符串"）
|
||||
将你认为最好的两个实现，放在最前面。
|
||||
比如：
|
||||
{my_poly_eval1, "超级垃圾实现"},
|
||||
{my_poly_eval2, "好一点的实现"},
|
||||
*/
|
||||
|
||||
peval_fun_rec peval_fun_tab[] =
|
||||
{
|
||||
|
||||
/* 第一项，应当是你写的最好CPE的函数实现 */
|
||||
{cuda_poly_eval_low_cpe, "CUDA optimized low CPE implementation"},
|
||||
/* 第二项，应当是你写的在10阶时具有最好性能的实现 */
|
||||
{cuda_poly_eval_degree10, "CUDA optimized degree 10 implementation"},
|
||||
|
||||
{poly_eval, "poly_eval: 参考实现"},
|
||||
|
||||
/* 下面的代码不能修改或者删除！！表明数组列表结束 */
|
||||
{NULL, ""}
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
Reference in New Issue
Block a user