【CUDA编程】GPU上的topK程序

原创 已于 2025-03-28 19:59:49 修改 · 266 阅读 · 0 ·
CC 4.0 BY-SA版权
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
文章标签：

#算法 #c++ #gpu算力 #深度学习 #CUDA

于 2025-03-21 10:47:14 首次发布

CUDA编程的一些例子

最近在学习CUDA编程，写了一点简单的程序，刚开始用学习GPU编程，踩了很多坑，特此记录一下

基于堆排序的topK

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <assert.h>
#include <stdio.h>
#include <stdlib.h>


#define K 10

__device__ void mySwap(int& num1, int& num2) {
	int tmp = num1;
	num1 = num2;
	num2 = tmp;
}

// 保持堆的特性
__device__ void heapify(int* data, int idx, int maxSize) {
	if (idx >= maxSize)
		return;
	int left = idx * 2 + 1, right = idx * 2 + 2;
	int maxIdx = idx;
	if (left < maxSize && data[left] < data[maxIdx])
		maxIdx = left;
	if (right < maxSize && data[right] < data[maxIdx])
		maxIdx = right;
	if (maxIdx != idx) {
		mySwap(data[idx], data[maxIdx]);
		heapify(data, maxIdx, maxSize);
	}
}


template<
int NUM_PER_BLOCK,
int NUM_PER_THREAD
>
__global__ void myTopK(int* input, int* output) {
	__shared__ int sdata[NUM_PER_BLOCK / NUM_PER_THREAD * K];
	int tid = threadIdx.x;
	int idx = blockDim.x * blockIdx.x + tid;

	int data[K];

	for (int i = K - 1; i >= 0; i--) {
		data[i] = input[i + NUM_PER_THREAD * idx];
		heapify(data, i, K);
	}
	for (int i = K; i < NUM_PER_THREAD; i++) {
		if (input[i + NUM_PER_THREAD * idx] > data[0]) {
			data[0] = input[i + NUM_PER_THREAD * idx];
			heapify(data, 0, K);
		}
	}
	for (int i = 0; i < K; i++) {
		sdata[idx * K + i] = data[i];
	}
	__syncthreads();
	int curNum = NUM_PER_BLOCK * K / NUM_PER_THREAD;

	while (true) {
		if (idx * NUM_PER_THREAD >= curNum)
			break;

		int sortNum = min((idx + 1) * NUM_PER_THREAD, curNum) - idx * NUM_PER_THREAD;
		if (sortNum <= K) {
			for (int i = 0; i < sortNum; i++) {
				data[i] = sdata[idx * NUM_PER_THREAD + i];
			}
		}
		else {
			for (int i = K - 1; i >= 0; i--) {
				data[i] = sdata[i + NUM_PER_THREAD * idx];
				heapify(data, i, K);
			}
			int end = min(NUM_PER_THREAD, sortNum); // 目前需要排序的数据量可能位于NUM_PER_THREAD和K之间，即NUM_PER_THREAD>sortNum>K
			for (int i = K; i < end; i++) {
				if (sdata[i + NUM_PER_THREAD * idx] > data[0]) {
					data[0] = sdata[i + NUM_PER_THREAD * idx];
					heapify(data, 0, K);
				}
			}
			sortNum = K;
		}
		__syncthreads();

		if (idx * NUM_PER_THREAD < curNum) {
			#pragma unroll
			for (int i = 0; i < sortNum; i++) {
				sdata[idx * K + i] = data[i];
			}
		}
		__syncthreads();

		curNum = curNum * K / NUM_PER_THREAD;
	}
	if (tid == 0) {
		for (int i = 0; i < K; i++) {
			output[i] = sdata[i];
		}
	}
}

int main() {
	const int BLOCK_SIZE = 256;
	const int N = 1024 * 1024;
	const int NUM_PER_BLOCK = N; // 1024
	const int NUM_PER_THREAD = NUM_PER_BLOCK / BLOCK_SIZE; // 4096

	assert(K < NUM_PER_THREAD);

	int* input = (int*)malloc(N * sizeof(int));
	int* output = (int*)malloc(K * sizeof(int));
	int* d_input;
	int* d_output;
	cudaError_t error = cudaMalloc(&d_input, N * sizeof(int));
	error = cudaMalloc(&d_output, K * sizeof(int));

	for (int i = 0; i < N; i++) {
		input[i] = rand() % 100;
	}

	for (int i = 0; i < K; i++) {
		int randIdx = rand() % N;
		input[randIdx] = i + 100;
		printf("input[%d] = %d\n", randIdx, input[randIdx]);
	}
	error = cudaMemcpy(d_input, input, N * sizeof(int), cudaMemcpyHostToDevice);

	dim3 BlockSize(BLOCK_SIZE, 1);
	myTopK <NUM_PER_BLOCK, NUM_PER_THREAD><< <1, BlockSize >> > (d_input, d_output);
	error = cudaMemcpy(output, d_output, K * sizeof(int), cudaMemcpyDeviceToHost);

	for (int i = 0; i < K; i++) {
		printf("output[%d] = %d.\n", i, output[i]);
	}
}

2路topK算法

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <assert.h>
#include <stdio.h>
#include <stdlib.h>



__device__ void mySwap(int& num1, int& num2) {
	int tmp = num1;
	num1 = num2;
	num2 = tmp;
}

// 保持堆的特性
__device__ void heapify(int* data, int idx, int maxSize) {
	if (idx >= maxSize)
		return;
	int left = idx * 2 + 1, right = idx * 2 + 2;
	int maxIdx = idx;
	if (left < maxSize && data[left] < data[maxIdx])
		maxIdx = left;
	if (right < maxSize && data[right] < data[maxIdx])
		maxIdx = right;
	if (maxIdx != idx) {
		mySwap(data[idx], data[maxIdx]);
		heapify(data, maxIdx, maxSize);
	}
}


template<
	int NUM_PER_BLOCK,
	int NUM_PER_THREAD,
	int K
>
__global__ void myTopK(int* input, int* output) {
	__shared__ int sdata[NUM_PER_BLOCK / NUM_PER_THREAD * K];
	int tid = threadIdx.x;
	int idx = blockDim.x * blockIdx.x + tid;
	if (idx * NUM_PER_THREAD >= NUM_PER_BLOCK) // 数据量较少时，后面的线程直接退出
		return;

	int data[K];

	for (int i = K - 1; i >= 0; i--) {
		data[i] = input[i + NUM_PER_THREAD * idx];
		heapify(data, i, K);
	}

	for (int i = K; i < NUM_PER_THREAD; i++) {
		if (input[i + NUM_PER_THREAD * idx] > data[0]) {
			data[0] = input[i + NUM_PER_THREAD * idx];
			heapify(data, 0, K);
		}
	}

	for (int i = 0; i < K; i++) {
		sdata[tid * K + i] = data[i];
	}
	__syncthreads();
	int curNum = NUM_PER_BLOCK * K / NUM_PER_THREAD;

	while (true) {
		if (tid * NUM_PER_THREAD >= curNum)
			break;

		int sortNum = min((tid + 1) * NUM_PER_THREAD, curNum) - tid * NUM_PER_THREAD;
		if (sortNum <= K) {
			for (int i = 0; i < sortNum; i++) {
				data[i] = sdata[tid * NUM_PER_THREAD + i];
			}
		}
		else {
			for (int i = K - 1; i >= 0; i--) {
				data[i] = sdata[i + NUM_PER_THREAD * tid];
				heapify(data, i, K);
			}
			int end = min(NUM_PER_THREAD, sortNum); // 目前需要排序的数据量可能位于NUM_PER_THREAD和K之间，即NUM_PER_THREAD>sortNum>K

			for (int i = K; i < end; i++) {
				if (sdata[i + NUM_PER_THREAD * tid] > data[0]) {
					data[0] = sdata[i + NUM_PER_THREAD * tid];
					heapify(data, 0, K);
				}
			}
			sortNum = K;
		}
		__syncthreads();
		if (tid * NUM_PER_THREAD < curNum) {

			for (int i = 0; i < sortNum; i++) {
				sdata[tid * K + i] = data[i];
			}
		}
		__syncthreads();

		curNum = curNum * K / NUM_PER_THREAD;
	}
	if (tid == 0) {
		for (int i = 0; i < K; i++) {
			output[blockIdx.x * K + i] = sdata[i];
		}
	}
}

int main() {
	const int BLOCK_SIZE = 256;
	const int GRID_SIZE = 16;
	const int N = 1024 * 1024;

	const int GLOBAL_K = 16;
	const int BLOCK_K = 32;

	assert(BLOCK_K % GLOBAL_K == 0);

	int* input = (int*)malloc(N * sizeof(int));
	int* output = (int*)malloc(GLOBAL_K * sizeof(int));
	int* d_input;
	int* d_output;
	int* d_block_output;
	cudaError_t error = cudaMalloc(&d_input, N * sizeof(int));
	error = cudaMalloc(&d_output, GLOBAL_K * sizeof(int));
	error = cudaMalloc(&d_block_output, BLOCK_K * GRID_SIZE * sizeof(int));
	for (int i = 0; i < N; i++) {
		input[i] = rand() % 100;
	}

	for (int i = 0; i < GLOBAL_K; i++) {
		int randIdx = rand() % N;
		input[randIdx] = i + 100;
		printf("input[%d] = %d\n", randIdx, input[randIdx]);
	}
	error = cudaMemcpy(d_input, input, N * sizeof(int), cudaMemcpyHostToDevice);

	dim3 GridSize(GRID_SIZE, 1);
	dim3 BlockSize(BLOCK_SIZE, 1);
	myTopK <N / GRID_SIZE, N / (GRID_SIZE * BLOCK_SIZE), BLOCK_K> << <GridSize, BlockSize >> > (d_input, d_block_output);

	assert(BLOCK_K * GRID_SIZE / GLOBAL_K > GLOBAL_K); // 确保每个线程处理的数据量大于K
	myTopK <BLOCK_K * GRID_SIZE, BLOCK_K * GRID_SIZE / GLOBAL_K, GLOBAL_K> << <1, BlockSize >> > (d_block_output, d_output);
	error = cudaMemcpy(output, d_output, GLOBAL_K * sizeof(int), cudaMemcpyDeviceToHost);

	for (int i = 0; i < GLOBAL_K; i++) {
		printf("output[%d] = %d.\n", i, output[i]);
	}
}
GPU编程与优化——OpenCL
AI天才研究院
09-03 1700
作者：禅与计算机程序设计艺术 1.简介 OpenCL（Open Computing Language）是一个开源的、跨平台、通用且功能丰富的编程模型，它提供一个框架，通过可移植编译器支持多种设备平台（CPU/GPU/FPGA），并通过统一接口使得开发者可以方便地部署在各种环境中运行，从而让程序在各个设备
CUDA专题3：为什么GPU能改变计算？深度剖析架构、CUDA®与可扩展编程
最新发布
03-29 456
图形处理器（GPU）在相近的成本和功耗范围内，能够提供比中央处理器（CPU）更高的指令吞吐量和内存带宽。许多应用程序利用这些优势，在 GPU 上的运行速度远超 CPU（参见《GPU 应用》）。其他计算设备（如 FPGA）虽然能效也很高，但其编程灵活性远不及 GPU。GPU 与 CPU 的能力差异源于它们的设计目标不同。CPU 旨在以最快速度执行单个线程（即一系列操作），并可并行执行数十个线程；而 GPU 则专为并行执行数千个线程优化（通过牺牲单线程性能来换取更高的整体吞吐量）。
[CUDA] Top_k问题
chiminmang7244的博客
05-02 1220
欢迎访问小站，阅读原文http://www.yandong.org/archives/759 Top_K问题是从一组数据中选出最大（最小）的K个数据。 问题分析 使用双调排序网络对数据进行排序，然后选择前k个数据。其中面临的问题是BLOCK之间的通信，可以使用global me...
CUDA编程模型系列九( topK 问题/规约/2_Pass核函数)
专注于人工智能领域的小何尚
07-06 2327
本期我介绍了cuda 当中规约算法的另一种情况, 利用跑两遍核函数来解决多个block中数据合并问题, 将第一遍的输出当做第二遍核函数的输入.Topk问题也是比较常用的方法, 希望我的解题思路能够帮到大家, 我是抛砖引玉, 希望大家能写出更快地核函数.本系列视频目的是帮助开发者们一步步地学会利用CUDA编程模型加速GPU应用, 我们的口号是: 让GPU飞起来。CUDA编程模型系列九( topK 问题/规约/2。
求topk——CUDA on ARM Platform夏令营 Day04 课后作业
weixin_48266461的博客
07-08 1246
本文利用CUDA求topk的思想是在第二部分向量元素求和思想的基础上得到的，读者理清楚后会更容易理解如何求topk。 有图解哦~ ヾ(✿ﾟ▽ﾟ)ノ
[CUDA 学习笔记] 百度 STI2 赛题二-基于向量交集的 TopK 搜索 优化学习笔记
LostUnravel的博客
12-18 734
给定850万条规模的数据文件，每条数据是最大128维度的整型id向量 （称为doc），id取值范围是0-50000，给定一个最大128维的整型id向量（称为query），求query与doc全集内各数据的交集个数topk（k=100）topk_v0_base.cu 为赛题的 Baseline 实现. 该实现整体思路比较简单, 首先是主机端内存分配等预处理, 然后在 GPU 上进行求交集, 最后在主机端根据求交集的分数降序排序得到最终结果.这个 Baseline 主要值得一提的是其 docs 的 swizzl
cuda编程（示例，topk问题，代码解读与cuda基础概念）
qq_43616651的博客
12-13 1087
使用锁页内存当数据传输频繁且对带宽要求高。使用零拷贝内存对于大量但访问频率较低的数据。使用托管内存以简化内存管理，尤其在数据访问模式不确定的情况下。每种内存类型都有其优势和局限，选择合适的类型取决于具体的应用需求和性能目标。好了，以上是一些基础概念。继续回到代码。总的思路即就是在每个block中求出最大的前20个值，放入中间结果数组，然后再执行一遍核函数对中间结果求前20个，最后输出结果。这样为什么如下定义应该就明白了。#define N 100000000 //数据大小。
GPU加速的分段Top-k查询算法.pdf
09-25
本文提出了一种基于CUDA（Compute Unified Device Architecture）模型的GPU加速分段Top-k查询算法。CUDA是NVIDIA公司开发的一种编程模型，允许开发者利用GPU的并行计算能力来解决复杂的计算问题。在分段查询算法中，...
cuda面试常见编程题目
m0_38086244的博客
10-11 757
1.写一个ReduceKernel。4.写一个矩阵旋转90°。
Qwen2-VL出现RuntimeError: CUDA error: too many resources requested for launch CUDA kernel errors的解决方案
热门推荐
weixin_43178406的博客
10-31 13万+
本文主要介绍了Qwen2-VL出现RuntimeError: CUDA error: too many resources requested for launch CUDA kernel errors的解决方案，希望能对使用Qwen2-VL的同学们有所帮助。 文章目录 1. 问题描述 2. 解决方案
top-k-pruned-landmark-labeling:Top-K PLL（AAAI '15）的实现
05-12
Top-K修剪地标标签 Top-K修剪地标标记是一种快速算法，用于回答现实网络（例如社交网络和网络图）上的top-k距离查询。 用法 从CUI界面 $ make $ bin/construct_index sample/example.txt 16 0 sample/index_file #compute an index that answer top 16 shortest distance on an undirected graph specified by sample/example.txt. $ bin/k_distance sample/index_file <<< "9 12" 16 1 2 3 3 3 3 3 3 3 3 4 4 4 4 4 4 执行make生成程序。 执行bin/construct_index以从给定的图形构造索引。 构造的索引存储在ind
CUDA入门之利用GPU寻找一组数据中最大的k个元素（一）
Not_answer的博客
07-09 2181
利用“插入排序”实现CUDA优化寻找一组数据中最大的k个元素
CUDA入门之利用GPU寻找一组数据中最大的k个元素（二）
Not_answer的博客
07-09 809
采用mul-pass的思路优化2-pass的算法，得到了较好的表现
topK问题（大顶堆、快速选择算法）
MRwangyd的博客
05-18 718
topK问题：有 N (N>1000000)个数,求出其中的前K个最小的数。 力扣原题：最小的k个数 输入整数数组 arr ，找出其中最小的 k 个数。 方法一：大顶堆 思路：维护一个大小为k的大顶堆，遍历一次数组，初始插入k个数，然后每遍历一个数，将其与堆顶比较，若比堆顶小，则堆顶弹出，该数入堆 class Solution { public: vector<int> getLeastNumbers(vector<int>& arr, int k) {
CUDA编程（二）
zhanglehes的专栏
08-28 628
概要 源于facebook的faiss库，其用过GPU 对于进行加速，另外xfr中也用到GPU，而需求源于之前xid对于大数据量的性能问题。之前也了解过一些机器学习的平台，如tersorflow，paddlepaddle，他们均将GPU加速作为平台的一大“卖点”。因此思考GPU编程对于智能基础服务还是有很大的帮助，于是“贸然”涉足了这一领域。 CUDA编程给我的感觉是 1. 入门很容易，我大概...
解决TOPK问题与堆排序
m0_53975470的博客
11-08 400
两种方法解决TopK问题和堆排序
cuda_knn vs torch.topk()
alfred_torres的博客
03-21 2551
发现一个快速的KNN gpu算法 KNN_CUDA DGCNN KNN 之前的knn算法用的都是DGCNN里面的torch.topk(),但是效率很低，网络训练既占用显存又慢；代替的pointnet2中的ball_query方法又不是严格的knn，可能会导致训练精度变低。 ### DGCNN knn def knn(x, k): inner = -2*torch.matmul(x.transpose(2, 1), x) xx = torch.sum(x**2, dim=1, keepdim=
【亲测免费】 GPU-TopK：加速GPU上的高效Top-K处理方案
gitblog_00935的博客
08-29 870
GPU-TopK：加速GPU上的高效Top-K处理方案 项目介绍 GPU-TopK是一个专为图形处理器（GPUs）设计的开源项目，旨在解决数据处理中的一项关键问题——从一组键值对或仅含键的数据数组中，高效地找到价值最大的前K个条目。利用GPU的强大并行计算能力，GPU-TopK提供了三种不同的实现策略：基于位旋排序的Bitonic Top-K、径向选择算法的Radix Select Top-K以及...
【亲测免费】 GPU-TopK：GPU上的高效Top-K查询处理库
gitblog_00006的博客
06-17 858
GPU-TopK：GPU上的高效Top-K查询处理库 项目简介 GPU-TopK是一个专为GPU设计的高效的Top-K运行时库。该库旨在解决一个基础问题：给定一个仅含键或键值对的数组，如何基于键的值找到排名最高的前k个元素？这个库提供了三种不同的算法实现，包括Bitonic Top-K、Radix Select Top-K和Sort Top-K。 项目技术分析 Bitonic Top-K：基于位...
unable to lock gpu to purge memory。invoked oom killer
03-16
如果系统无法锁定GPU内存进行清理，可能是因为GPU驱动程序存在问题，或者有进程正在占用GPU资源，导致无法释放内存。这时候，系统可能会触发OOM Killer（Out-Of-Memory Killer）来终止进程以释放内存。 接下来，我...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值