使用一个一维数组进行矩阵转置以及其并行化

3 篇文章 0 订阅

直接上代码,思路全在纸上,写两下找规律就能写了

#include<time.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<stdlib.h>
#include <stdio.h>
#define maxsize 2000
#define INDEX(i, j, ldy)  ((i)* (ldy)+(j))

__global__ void Turn(int* data, int n ){
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	for (int j = i + 1; j < n - i; j++) {
		int temp = 0;
		temp = data[INDEX(i, j, maxsize)];//保存了左上角的第一个点
		data[INDEX(i, j, maxsize)] = data[INDEX(j, maxsize - 1 - i, maxsize)];//左上角的点=右上角的点
		data[INDEX(j, maxsize - 1 - i, maxsize)] = data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)];//右上角的点=右下角的点
		data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)] = data[INDEX(maxsize - 1 - j, i, maxsize)];//右下角的点=左下角的点
		data[INDEX(maxsize - 1 - j, i, maxsize)] = temp;//左下角的点=左上角的点

		//向右旋转:
		//i!j=n-1-j !i
		//n-1-j!i=n-1-i!n-1-j
		//n-1-i!n-1-j=n-1-i
		//j!n-1-i=temp
	}
}
void Print(int* data) {
	for (int i = 0; i < maxsize; i++)
	{
		for (int j = 0; j < maxsize; j++)
		{
			printf(" %d", data[INDEX(i, j, maxsize)]);

		}
		printf("%\n");
	}
	printf("%\n");
	printf("%\n");
}
int* CPU(int* data, int n, int x) {
	//利用循环判断每次该旋转几个数,像剥洋葱一样从外到里一层一层操作
	//每次都旋转一层,每层从i+1开始到n-i-1为止
	//每次旋转四个对角的数
	for (int i = 0; i < n / 2; i++) {
		for (int j = i + 1; j < n - i; j++) {
			int temp = 0;
			temp = data[INDEX(i, j, maxsize)];//保存了左上角的第一个点
			data[INDEX(i, j, maxsize)] = data[INDEX(j, maxsize - 1 - i, maxsize)];//左上角的点=右上角的点
			data[INDEX(j, maxsize - 1 - i, maxsize)] = data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)];//右上角的点=右下角的点
			data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)] = data[INDEX(maxsize - 1 - j, i, maxsize)];//右下角的点=左下角的点
			data[INDEX(maxsize - 1 - j, i, maxsize)] = temp;//左下角的点=左上角的点

			//向右旋转:
			//i!j=n-1-j !i
			//n-1-j!i=n-1-i!n-1-j
			//n-1-i!n-1-j=n-1-i
			//j!n-1-i=temp
		}
	}
	//Print(data);
	return data;
}
int Check(int* a, int* b) {
	for (int i = 0; i < maxsize; i++)
	{
		for (int j = 0; j < maxsize; j++)
		{
			if (a[INDEX(i, j, maxsize)] != b[INDEX(i, j, maxsize)])
				return 1;
		}
	}
	return 0;
}
void GPU(int* data) {


	
	printf("\n	GPU code is running...\n");
	int n = maxsize;
	int* d_data;
	cudaMalloc((void**)&d_data, sizeof(int) * maxsize * maxsize);
	cudaMemcpy(d_data, data, sizeof(int) * maxsize * maxsize, cudaMemcpyHostToDevice);
	int block_num = maxsize/2;
	int grid_num = (maxsize + block_num - 1) / block_num;


	clock_t start1, end1;
	start1 = clock();
	Turn << < grid_num, block_num >> > (d_data, n);
	end1 = clock();
	float time1 = (float)(end1 - start1) / CLOCKS_PER_SEC;
	printf("	The tim GPU used is:%lf\n", time1);
	cudaMemcpy(data,d_data, sizeof(int) * maxsize * maxsize, cudaMemcpyDeviceToHost);

	printf("	GPU program is end...\n\n");
	cudaFree(d_data);
}
int main() {

	int* a;
	cudaMallocHost((void**)&a, sizeof(int) * maxsize * maxsize);
	for (int i = 0; i < maxsize; i++)
	{
		for (int j = 0; j < maxsize; j++)
		{
			a[INDEX(i, j, maxsize)] = INDEX(i, j, maxsize);
		}
	}
	int* b;
	b = (int*)malloc(sizeof(int) * maxsize * maxsize);
	cudaMemcpy(b, a, sizeof(int) * maxsize * maxsize, cudaMemcpyHostToHost);

	//printf("数组创建成功:\n");
	//Print(a);


	printf("\n	CPU code is running...\n");
	clock_t start1, end1;
	start1 = clock();
	a = CPU(a, maxsize, 0);
	end1 = clock();
	float time1 = (float)(end1 - start1) / CLOCKS_PER_SEC;
	printf("	The tim CPU used is:%lf\n",time1);
	printf("	CPU program is end...\n\n");

	//printf("数组b:\n");
	//Print(b);
	//printf("数组a旋转后:\n");
	//Print(a);

	GPU(b);
	if (Check(a, b)) {
		printf("\nthe GPU's answer is wrong!\n");
	}
	else {
		printf("	The Answer is Right!\n");
	}
	free(b);
	cudaFree(a);

	return 0;
}

目前看起来并行化的加速比性能很高,但是我把最大数参数maxsize调到2000以上的时候,答案就开始错误了,不知道为什么,后面解决后再更新,欢迎大家一起探讨。


4月6日更新


通过对每一个过程采用
cudaStatus != cudaSuccess的方式,一步步排查出了问题,发现原来是我的最大线程数设置了为最大数/2,而我电脑支持的最大线程数却只有1024,超过之后就会报参数错误:invalid configuration argument

修改后的代码已经可以正确运行并且没有出错了。

#include<time.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<stdlib.h>
#include <stdio.h>
#define maxsize 10000
#define INDEX(i, j, ldy)  ((i)* (ldy)+(j))

/*
小写一个思路:
写4个GPU函数,每一个对应矩阵的1/4的数量
但是我是用一维数组来保存的展开的二维矩阵,所以需要设计如何保存这1/4

*/
__global__ void Turn(int* data, int n ){
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	for (int j = i + 1; j < n - i; j++) {
		int temp = 0;
		temp = data[INDEX(i, j, maxsize)];//保存了左上角的第一个点
		data[INDEX(i, j, maxsize)] = data[INDEX(j, maxsize - 1 - i, maxsize)];//左上角的点=右上角的点
		data[INDEX(j, maxsize - 1 - i, maxsize)] = data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)];//右上角的点=右下角的点
		data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)] = data[INDEX(maxsize - 1 - j, i, maxsize)];//右下角的点=左下角的点
		data[INDEX(maxsize - 1 - j, i, maxsize)] = temp;//左下角的点=左上角的点

		//向右旋转:
		//i!j=n-1-j !i
		//n-1-j!i=n-1-i!n-1-j
		//n-1-i!n-1-j=n-1-i
		//j!n-1-i=temp
	}
}
void Print(int* data) {
	for (int i = 0; i < maxsize; i++)
	{
		for (int j = 0; j < maxsize; j++)
		{
			printf(" %d", data[INDEX(i, j, maxsize)]);

		}
		printf("%\n");
	}
	printf("%\n");
	printf("%\n");
}
int* CPU(int* data, int n, int x) {
	//利用循环判断每次该旋转几个数,像剥洋葱一样从外到里一层一层操作
	//每次都旋转一层,每层从i+1开始到n-i-1为止
	//每次旋转四个对角的数
	for (int i = 0; i < n / 2; i++) {
		for (int j = i + 1; j < n - i; j++) {
			int temp = 0;
			temp = data[INDEX(i, j, maxsize)];//保存了左上角的第一个点
			data[INDEX(i, j, maxsize)] = data[INDEX(j, maxsize - 1 - i, maxsize)];//左上角的点=右上角的点
			data[INDEX(j, maxsize - 1 - i, maxsize)] = data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)];//右上角的点=右下角的点
			data[INDEX(maxsize - 1 - i, maxsize - 1 - j, maxsize)] = data[INDEX(maxsize - 1 - j, i, maxsize)];//右下角的点=左下角的点
			data[INDEX(maxsize - 1 - j, i, maxsize)] = temp;//左下角的点=左上角的点

			//向右旋转:
			//i!j=n-1-j !i
			//n-1-j!i=n-1-i!n-1-j
			//n-1-i!n-1-j=n-1-i
			//j!n-1-i=temp
		}
	}
	//Print(data);
	return data;
}
int Check(int* a, int* b) {
	for (int i = 0; i < maxsize; i++)
	{
		for (int j = 0; j < maxsize; j++)
		{
			if (a[INDEX(i, j, maxsize)] != b[INDEX(i, j, maxsize)])
				return 1;
		}
	}
	return 0;
}
cudaError_t GPU(int* data) {
	printf("\n	GPU code is running...\n");
	int n = maxsize;
	int* d_data;
	cudaError_t cudaStatus;
	cudaStatus =cudaMalloc((void**)&d_data, sizeof(int) * maxsize * maxsize);
	if (cudaStatus != cudaSuccess) {
		printf("device memory malloc allocation failed"); 
		return cudaStatus;
	}
	cudaStatus = cudaMemcpy(d_data, data, sizeof(int) * maxsize * maxsize, cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		printf("device memory allocation failed");
		return cudaStatus;
	}
	
	int block_num = 1024;
	int grid_num = (maxsize + block_num - 1) / block_num;

	clock_t start1, end1;
	start1 = clock();
	Turn << < grid_num, block_num >> > (d_data, n);
	end1 = clock();
	cudaStatus = cudaGetLastError();
		if (cudaStatus != cudaSuccess) 
		{
			fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
			return cudaStatus;
		}
		cudaStatus = cudaDeviceSynchronize();
		 if (cudaStatus != cudaSuccess) {
			 printf("Synchronize failed!\n");
			 return cudaStatus;
		 }

	float time1 = (float)(end1 - start1) / CLOCKS_PER_SEC;
	printf("	The tim GPU used is:%lf\n", time1);
	cudaStatus = cudaMemcpy(data, d_data, sizeof(int) * maxsize * maxsize, cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess) {
		printf("Device to Host memory  failed");
		return cudaStatus;
	}
	

	printf("	GPU program is end...\n\n");
	cudaFree(d_data);
	return cudaSuccess;
}
int main(void ) {

	int* a;
	cudaMallocHost((void**)&a, sizeof(int) * maxsize * maxsize);
	for (int i = 0; i < maxsize; i++)
	{
		for (int j = 0; j < maxsize; j++)
		{
			a[INDEX(i, j, maxsize)] = INDEX(i, j, maxsize);
		}
	}
	int* b;
	b = (int*)malloc(sizeof(int) * maxsize * maxsize);
	cudaMemcpy(b, a, sizeof(int) * maxsize * maxsize, cudaMemcpyHostToHost);

	//printf("数组创建成功:\n");
	//Print(a);
	printf("\n	CPU code is running...\n");
	clock_t start1, end1;
	start1 = clock();
	a = CPU(a, maxsize, 0);
	end1 = clock();
	float time1 = (float)(end1 - start1) / CLOCKS_PER_SEC;

	printf("	The tim CPU used is:%lf\n",time1);
	printf("	CPU program is end...\n\n");

	//printf("数组b:\n");
	//Print(b);
	//printf("数组a旋转后:\n");
	//Print(a);
	cudaError_t cudaStatus;
	cudaStatus=GPU(b);
	if (cudaStatus != cudaSuccess) {
		printf("GPU failed!");
		return 1;
	}
	if (Check(a, b)) {
		printf("\nthe GPU's answer is wrong!\n");
	}
	else {
		printf("	The Answer is Right!\n");
	}
	free(b);
	cudaFree(a);

	return 0;
}
  • 0
    点赞
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值