现在排序算法都是基于CPU的,但是CPU有它的并发局限性
GPU需要处理大量三角形数据和渲染数据,并行计算能力非常强,那么为什么不可以利用GPU的超强并行能力完成排序
(笔者曾经看过FPGA编程里面的排序算法,直接利用大规模并行电路直接完成排序,GPU和FAPG居然有一些相似)
为了追求极致的效率为什么不可以把显卡调用起来取完成排序呢
想到就去做,下载cuda,可以对cuda完成面向显卡的编程,按教程安装好环境好,发现运行失败,后面发现是显卡算力的问题,查了资料以后解决了
VS2022搭建cuda环境的no kernel image is available for execution on the device问题解决_星空_MAX的博客-CSDN博客
环境为VS2022和cuda11.6
原理
比如给出5个无序数字
4,3,1,2,5
开启五个并行线程
分别统计每个数字比它小的数字有几个
4比较完4,3,1,2,5发现小于等于的数字有4个,那么把新数组的第4个位置放置为4
这是的新数组为:
空 空 空 4 空
3比较完4,3,1,2,5发现比自己小的数字有3个,那么把新数组的第3个位置放置为3
这是的新数组为:
空 空 3 4 空
五个数字分别在五个并行线程中找到了自己的位置:
最后的数组为:
1,2,3,4,5
以上就完成了秩排序
代码
#include <stdio.h>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>
#ifndef __CUDACC__
#define __CUDACC__
#endif
#define ARRAY_SIZE 6
#define THREADS_PER_BLOCK 6
__global__ void rankSort(int* d_a, int* d_b)
{
int count = 0;
int tid = threadIdx.x;
int ttid = threadIdx.x + blockIdx.x * blockDim.x;
int val = d_a[ttid];
__shared__ int cache[THREADS_PER_BLOCK];
for (int i = tid; i < ARRAY_SIZE; i += THREADS_PER_BLOCK)
{
cache[tid] = d_a[i];
//等待当前块所有线程完成共享内存数据填充
__syncthreads();
//统计共享内存中小于当前值得数量
for (int j = 0; j < THREADS_PER_BLOCK; j++)
{
if (val > cache[j])
{
count++;
//等待所有线程完成统计工作
__syncthreads();
}
}
}
d_b[count] = val;
}
int main()
{
//定义主机(CPU)和设备(GPU)变量
int h_a[ARRAY_SIZE] = { 5, 9, 3, 4, 1, 8 };
int h_b[ARRAY_SIZE];
int* d_a, * d_b;
//分配设备内存
cudaMalloc(&d_a, ARRAY_SIZE * sizeof(int));
cudaMalloc(&d_b, ARRAY_SIZE * sizeof(int));
//拷贝数据:从CPU到GPU
cudaMemcpy(d_a, h_a, ARRAY_SIZE * sizeof(int), cudaMemcpyHostToDevice);
//调用内核函数
rankSort << <ARRAY_SIZE / THREADS_PER_BLOCK, THREADS_PER_BLOCK >> > (d_a, d_b); //从这里传入两个数组
//等待内核函数运行完成
cudaDeviceSynchronize();
//拷贝数据:从GPU到CPU
cudaMemcpy(h_b, d_b, ARRAY_SIZE * sizeof(int), cudaMemcpyDeviceToHost);
//打印秩排序结果
printf("result:");
for (int i = 0; i < ARRAY_SIZE; i++)
{
printf("%d ", h_b[i]);
}
//释放GPU内存
cudaFree(d_a);
cudaFree(d_b);
system("pause");
return 0;
}
可以发现排序成功了