文章目录
VS2017 CUDA编程学习1:CUDA编程两变量加法运算
VS2017 CUDA编程学习2:在GPU上执行线程
VS2017 CUDA编程学习3:CUDA获取设备上属性信息
VS2017 CUDA编程学习4:CUDA并行处理初探 - 向量加法实现
VS2017 CUDA编程学习5:CUDA并行执行-线程
VS2017 CUDA编程学习6: GPU存储器架构
VS2017 CUDA编程学习7:线程同步-共享内存
VS2017 CUDA编程学习8:线程同步-原子操作
VS2017 CUDA编程学习9:常量内存
VS2017 CUDA编程学习10:纹理内存
VS2017 CUDA编程学习实例1:CUDA实现向量点乘
VS2017 CUDA编程学习11:CUDA性能测量
VS2017 CUDA编程学习12:CUDA流
前言
今天跟大家分享如何使用CUDA实现秩排序。
1. C++ CUDA实现秩排序
这里使用了共享内存和同步的概念实现秩排序
详细代码如下所示:
#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 5
#define THREADS_PER_BLOCK 5
__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, 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("秩排序的结果:\n");
for (int i = 0; i < ARRAY_SIZE; i++)
{
printf("%d\n", h_b[i]);
}
//释放GPU内存
cudaFree(d_a);
cudaFree(d_b);
system("pause");
return 0;
}
3. 执行结果
总结
为了更好的理解CUDA编程的理论知识,实践是必不可少的,只有更多的实践才能更好的理解新的知识,实践出真知啊。
学习资料
《基于GPU加速的计算机视觉编程》