直接上代码,思路全在纸上,写两下找规律就能写了
#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;
}