cuda笔记-GPU多线程的奇偶排序

偏执的太偏执、 2022-12-28 07:39 192阅读 0赞

首先说明奇偶排序：

算法的思路是先排奇数序号的相邻2个，或者偶数序号的相邻两个，然后一直到序列有序为止，如下代码：

#include<stdio.h>
    #include<stdlib.h>
    #include<stdbool.h>
    #include <iostream>
    
    using namespace std;
    
    
    void swap(int *a, int *b)
    {
    	int t;
    	t = *a;
    	*a = *b;
    	*b = t;
    }
    void printArray(int a[], int count)
    {
    	int i;
    	for (i = 0; i < count; i++)
    		printf("%d ", a[i]);
    	printf("\n");
    }
    void Odd_even_sort(int a[], int size)
    {
    	bool sorted = false;
    	int count = 1;
    	while (!sorted)
    	{
    		sorted = true;
    		for (int i = 1; i < size - 1; i += 2)
    		{
    			if (a[i] > a[i + 1])
    			{
    				swap(&a[i], &a[i + 1]);
    				sorted = false;
    			}
    		}
    		for (int i = 0; i < size - 1; i += 2)
    		{
    			if (a[i] > a[i + 1])
    			{
    				swap(&a[i], &a[i + 1]);
    				sorted = false;
    			}
    		}
    		
    		cout << "count:" << count << endl;
    	}
    }
    int main(void)
    {
    	int a[] = { 9, 8, 7, 6, 5, 4, 3, 2, 1 };
    	int n = sizeof(a) / sizeof(*a);
    	Odd_even_sort(a, n);
    	printArray(a, n);
    	getchar();
    	return 0;
    }

程序运行截图如下：

![watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3FxNzg0NDI3NjE_size_16_color_FFFFFF_t_70][]

这个奇偶排序还有个特点：

奇偶排序实际上在多处理器环境中很有用，处理器可以分别同时处理每一个奇数对，然后又同时处理偶数对。因为奇数对是彼此独立的，每一刻都可以用不同的处理器比较和交换。这样可以非常快速地排序。

下面就用cuda来实现下，其中有个要说明的地方：

cuda中\_\_syncthreads()的作用：block内部用于线程同步，就是同一block内所有线程执行至\_\_syncthreads()处等待全部线程执行完毕后再继续。

代码如下：

#include "cuda_runtime.h"
    #include "device_launch_parameters.h"
    #include <iostream>
    #include "curand.h"
    #include "curand_kernel.h"
    #include <stdio.h>
    #include <iostream>
    
    using namespace std;
    
    #pragma comment(lib, "cudart.lib")
    #pragma comment(lib, "curand.lib")
    
    void Matrix_init(float *a, int N) {
    
    	curandGenerator_t gen;
    	curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_MRG32K3A);
    	curandSetPseudoRandomGeneratorSeed(gen, 11ULL);
    	curandGenerateUniform(gen, a, N);
    }
    
    
    
    
    __global__ void sort(float *a, int N) {
    
    	int x = threadIdx.x;
    	extern __shared__ float s_a[];
    	s_a[x] = a[x];
    	__syncthreads();
    
    	for (int i = 0; i < N; i++)
    	{
    		int j = i % 2;
    		int idx = 2 * x + j;
    		if (idx + 1 < N && s_a[idx] < s_a[idx + 1]) {
    			float tmp = s_a[idx];
    			s_a[idx] = s_a[idx + 1];
    			s_a[idx + 1] = tmp;
    		}
    		__syncthreads();
    	}
    
    	a[x] = s_a[x];
    	__syncthreads();
    }
    
    int main() {
    
    	int m = 16;
    	int N = 256;
    
    	
    	float *p_d, *p_h, *p_hs;
    	p_h = (float*)malloc(N * sizeof(float));
    	p_hs = (float*)malloc(N * sizeof(float));
    	cudaMalloc((void**)&p_d, N * sizeof(float));
    
    	Matrix_init(p_d, N);
    	cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);
    	
    	for (int i = 0; i < N; i++) {
    
    		if (i % m == 0) {
    
    			printf("\n");
    		}
    		cout << " " << p_h[i] << " ";
    	}
    
    	cout << endl << "sort:" << endl;
    	sort << <1, N, N * sizeof(float) >> > (p_d, N);
    	cudaMemcpy(p_h, p_d, N * sizeof(float), cudaMemcpyDeviceToHost);
    	for (int i = 0; i < N; i++) {
    
    		if (i%m == 0) 
    			printf("\n");
    
    		cout << " " << p_h[i] << " ";
    	}
    
    
    	cudaFree(p_d);
    	free(p_h);
    	free(p_hs);
    
    	
    	getchar();
    	return 0;
    }

程序运行截图如下：

![watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3FxNzg0NDI3NjE_size_16_color_FFFFFF_t_70 1][]

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3FxNzg0NDI3NjE_size_16_color_FFFFFF_t_70]: /images/20221120/01d315edcc64426d83d8656e3fc98a81.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3FxNzg0NDI3NjE_size_16_color_FFFFFF_t_70 1]: /images/20221120/b52f89954733490ba392f887c59ae579.png