cuda编程（示例，topk问题，代码解读与cuda基础概念）

文章声明：本文代码来自于 扫地的小和尚，CUDA编程模型系列九( topK 问题/规约/2_Pass核函数)_cuda topk-CSDN博客

这篇文章，以及b站的视频 

CUDA编程模型系列九( topK 问题/规约/2_Pass核函数)_哔哩哔哩_bilibili

 可以自行观看，阅读原作，最近在学习如何使用cuda写程序，应用之前所学习的cuda基础知识，如何设计代码，如何编程，怎样去写一段代码是本文以及本人想要学习的。

1.topk问题

首先，先看看topk问题是什么，一个长度为N的数组（这里数据为整型int），求该数组前20个最大数分别是多少。

总结，输入的数组，长度为N，其次是输出数组，长度为k。

首先想想传统方法怎么做，遍历输入数组，使用排序或者比较算法先求最大值，存入数组，标记最大值的位置表示已访问，然后再求次最大值，标记已访问，再继续，直到数组存满。大循环是k次，小循环是N次，或者交换也行（就变成了插入排序，详情请看10中排序算法代码实现（Python与C++）-CSDN博客或者下面的代码）。这样使用两个for循环暴力求解问题，其实也能解决问题，就是对于海量的数据来说来耗时。

那么怎么用GPU去优化这个问题呢，首先要理解CPU的运行时逻辑运行，即一步一步走，一次循环结束，下一次循环才能开始，而GPU是一种并行运行，最小运行单位为线程，GPU中有大量的线程可以同步做运算执行一些任务。

对于上面的问题，在海量的数据中找出前20个最大值，那么用cuda来做可行的方法就是，将数据分成好多块由线程分别计算每个块里的前20个最大值，然后再合并到一个块里计算这个块里的前20个最大值。

这里要提一嘴，之前面试问我并行任务的前置条件是什么，我直接懵逼，后来查了以下才明白一个任务如果可以并行，那么这个任务就是可分割的。所以这一点是并行的思想，将一个问题分成几块或者好多块来解决这才是并行任务的前提。只能说非常后悔，但是没办法，错过了就是重来，就和谈恋爱似的。一般出现for循环的时候就可以考虑用并行来实现了。

好的回归正题，下面一边讲小和尚的cuda编程代码思路，一边复习下看用到了什么CUDA编程思想。

2.代码思路与代码段

首先，在全局领域定义了接下要用的常量以及变量。

#define N 100000000  //数据大小
#define BLOCK_SIZE 256  //一个块中有256个线程
#define GRID_SIZE 32  //32 每个网格中有32个块
#define topk 20 //topk问题前topk个最大数

__managed__ int source[N];   //原数组 
//__managed__  cuda关键字，用于声明所谓的托管内存，允许内存在CPU和GPU之间自动共享。
//用 __managed__ 声明的变量可以同时被 CPU 和 GPU 访问，无需手动在主机（CPU）和设备（GPU）之间复制数据。
//使用托管内存简化了内存管理，因为它允许 CPU 和 GPU 在无需显式数据传输命令的情况下访问相同的内存。

__managed__ int gpu_result[topk];  //topk最终结果
__managed__ int _1_pass_result[topk * GRID_SIZE];//每个block的前20个，即中间结果

一些代码解释也已经给出。这里分析下对应到那些cuda知识点。

由于主机端host（/cpu）和设备端device（GPU）端的内存是存在分离的。这意味着主机端不能直接访问设备端内存，而设备端也无法直接访问主机端内存。为了在两个环境之间交换数据，一般使用cudaMalloc（在GPU上分配内存），cudaMemcpy（将CPU内存数据复制到GPU内存/将GPU内存数据复制到CPU内存）最后使用cudaFree释放GPU上的内存。这也是cuda 提供的最常用，最基础的CPU，GPU交互方式。

但除此之外，还存在几种优化主机端与设备端的数据传输方式。

1.使用锁页内存/（固定内存），可以使用在cudaMallocHost在主机端申请锁页内存（用于存放主机端数据，使用cudaFreeHost（），清理锁页内存，设备端还是需要cudaMalloc申请），然后使用上面的基础方式完成数据交互。优点，这些内存时页面锁定的并且对设备来说可访问的，由于固定内存能被设备直接访问，所以它能用比可分页内存高得多的带宽进行读写。缺点，锁页内存有限，使用太多降低主机系统性能。还是需要显示地将数据从主机端存取内存数据。

2.零拷贝内存，常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机变量。但有一个例外，零拷贝内存。主机和设备都可以访问零拷贝内存。GPU设备可以直接访问零拷贝内存，而无需先将数据拷贝到GPU的内存中。零拷贝的使用，还是用cudaHostAlloc在主机端申请内存，在GPU使用时，只需要将在主机端申请的指针交给设备变量即可，使用cudaHostGetDevicePointer（）函数，最后在使用cudaFreeHost释放零拷贝内存。所以这里可以看到零拷贝就是将零拷贝内存的指针给设备变量，那么设备端就可直接访问了。

3.统一内存寻址与托管内存：统一内存寻址是一种内存架构，它允许CPU和GPU共享同一个地址空间，这意味着一个内存地址可以同时被CPU和GPU访问。这简化了数据传输和同步的复杂性，因为不需要专门的API调用来在CPU和GPU间复制数据。但是，需要注意的是，即使地址空间是统一的，实际的物理内存可能仍然是分离的，因此数据可能需要在CPU和GPU之间移动。

托管内存允许开发者分配内存，这些内存既可以被CPU访问，也可以被GPU访问，而无需担心数据在不同内存空间之间的显式复制。

在CUDA中使用托管内存的步骤通常如下：

1. 分配托管内存： 使用cudaMallocManaged分配托管内存。

2. 在CPU和GPU上使用内存： 既可以在CPU代码中直接访问这块内存，也可以在CUDA内核中访问。

3. 同步操作： 在CPU访问托管内存中的数据之前，使用cudaDeviceSynchronize确保所有GPU操作已经完成。

4. 释放内存： 使用cudaFree释放托管内存。

 所以第3种方式时最方便的，给数据申请一块内存，也不需要cudaMemcpy传来传去。

而在本文的代码中，使用了cuda编程中的__managed__关键字，用于声明托管内存（Managed Memory），这是一种简化主机和设备间内存交互的方式。当你在CUDA程序中使用 __managed__ 关键字声明变量时，这些变量会自动成为托管内存，即它们既可以被CPU访问，也可以被GPU访问，而无需进行显式的数据传输。当你在变量前使用 __managed__ 关键字时，CUDA运行时会负责这些变量的内存分配和释放，以及必要时在主机和设备之间自动移动数据。这样，你就可以在主机代码和设备代码（即CUDA内核）之间共享数据，而无需手动同步或传输数据。除此之外，在访问托管内存的数据之前，可能需要调用 cudaDeviceSynchronize() 来确保所有设备上的操作都已完成。

以上3种方式加上最基础的，一共四种，他们存在都是合理的，适合用于不同场景，虽然托管内存方便，但其性能收到PCIe总线带宽的限制，当频繁在主机和设备端交换数据时，并不理想。所以使用是要看场景的。

以下是由GPT总结的几种内存的使用场景。

锁页内存（Pinned Memory）

1. 适用场景:

   • 频繁数据传输：当需要频繁地在主机（CPU）和设备（GPU）之间传输数据时，使用锁页内存可以提高传输效率。
   • 高带宽要求：如果应用程序需要大量数据在CPU和GPU之间快速移动，锁页内存可以提供更高的数据传输带宽。

2. 不适用场景:

   • 内存受限：锁页内存不能被操作系统交换出去，因此在内存受限的情况下可能会影响系统性能。
   • 小规模数据传输：对于小量数据或不频繁的数据传输，使用锁页内存的优势不明显。

零拷贝内存（Zero-Copy Memory）

1. 适用场景:

   • 大量数据但低频访问：当处理的数据量非常大而且不经常访问时，零拷贝内存可以减少数据复制的开销。
   • 只读或写一次数据：对于只读一次或只写一次的数据，零拷贝内存能够减少不必要的数据拷贝。

2. 不适用场景:

   • 高频率数据访问：由于访问零拷贝内存的延迟较高，频繁访问这类内存会降低性能。
   • 小数据集：对于较小的数据集，传统的内存拷贝可能更高效。

托管内存（Managed Memory）

1. 适用场景:

   • 简化内存管理：当需要简化CPU和GPU间的内存管理时，托管内存提供了一种透明的数据迁移方式。
   • 动态数据迁移：适用于数据访问模式不易预测的场景，CUDA运行时会根据需要动态迁移数据。

2. 不适用场景:

   • 极致性能需求：如果需要精细控制内存以达到最优性能，手动管理内存可能更合适。
   • 确定的数据访问模式：如果数据访问模式非常明确，通过手动优化内存传输可能获得更好的性能。

总结

• 使用 锁页内存 当数据传输频繁且对带宽要求高。
• 使用 零拷贝内存 对于大量但访问频率较低的数据。
• 使用 托管内存 以简化内存管理，尤其在数据访问模式不确定的情况下。

每种内存类型都有其优势和局限，选择合适的类型取决于具体的应用需求和性能目标。

好了，以上是一些基础概念。继续回到代码。总的思路即就是在每个block中求出最大的前20个值，放入中间结果数组，然后再执行一遍核函数对中间结果求前20个，最后输出结果。这样为什么如下定义应该就明白了。

#define N 100000000  //数据大小
#define BLOCK_SIZE 256  //一个块中有256个线程
#define GRID_SIZE 32  //32 每个网格中有32个块
#define topk 20 

__managed__ int source[N];   //原数组 
//__managed__  cuda关键字，用于声明所谓的托管内存，允许内存在CPU和GPU之间自动共享。
//用 __managed__ 声明的变量可以同时被 CPU 和 GPU 访问，无需手动在主机（CPU）和设备（GPU）之间复制数据。
//使用托管内存简化了内存管理，因为它允许 CPU 和 GPU 在无需显式数据传输命令的情况下访问相同的内存。

__managed__ int gpu_result[topk];  //topk最终结果
__managed__ int _1_pass_result[topk * GRID_SIZE];//每个block的前20个，即中间结果

//理论，求一个大数组的前20个最大值，先将数组放入GPU内，每个block中求出最大的前20个值，放入_1_passresult
//然后每个block前20个值放一块在求前20个值得到最中结果

接下来就是main（）函数：

int main(){

	//为原数组赋初值
	printf("初始化源数据.....\n");
	for (int i = 0; i < N; i++) {
		source[i] = rand();
	}
	printf("完成初始化源数据.....\n");

	//cuda事件-计时
	cudaEvent_t start, stop_gpu, stop_cpu;
	cudaEventCreate(&start);
	cudaEventCreate(&stop_gpu);
	cudaEventCreate(&stop_cpu);
	cudaEventRecord(start);
	cudaEventSynchronize(start);//事件同步
	//这个函数用于等待一个 CUDA 事件完成。
	// 当你在 CUDA 程序中设置一个事件时，比如 cudaEventRecord(event, stream)，
	// 它会在特定的流（stream）中标记一个点。cudaEventSynchronize(event) 会阻塞调用线程，
	// 直到该事件发生，即直到 GPU 上的相关操作完成。
	printf("GPU Run *************\n");
	int times = 20;
	//计算
	for (int i = 0; i < times; i++) {
		gpu_topk << <GRID_SIZE, BLOCK_SIZE >> > (source, _1_pass_result, N, topk);
		gpu_topk << <1, BLOCK_SIZE >> > (_1_pass_result, gpu_result, topk * GRID_SIZE, topk);
		cudaDeviceSynchronize();
		//cudaDeviceSynchronize() 函数会阻塞调用线程，直到 GPU 完成所有队列中的操作。
		// 这包括所有 CUDA 核心、内存复制和其他相关的 GPU 操作。
	}
	printf("GPU Run Complete %d 次*************\n",times);
	cudaEventRecord(stop_gpu);
	cudaEventSynchronize(stop_gpu);


	//cpu结果初始化
	int cpu_result[topk] = { 0 }; //cpu结果存储
	printf("CPU Run *************\n");
	//计算
	cpu_topk(source, cpu_result, N, topk);
	printf("GPU Run Complete *************\n");
	cudaEventRecord(stop_cpu);
	cudaEventSynchronize(stop_cpu);

	//计算两次时间
	float time_cpu, time_gpu;
	cudaEventElapsedTime(&time_gpu, start, stop_gpu);
	cudaEventElapsedTime(&time_cpu, stop_gpu, stop_cpu);

	//判断GPU计算是否有误
	bool error = false;
	for (int i = 0; i < topk; i++) {
		printf(" CPU top%d: %d; Gputop%d: %d;\n", i + 1, cpu_result[i], i + 1, gpu_result[i]);
		if (fabs(gpu_result[i] - cpu_result[i]) > 0) {
			error = true;
		}
	}
	printf("Result:%s\n", (error ? "Error" : "pass"));
	printf("CPU time: %.2f; GPU time: %.2f\n", time_cpu, time_gpu);

	return 0;
}

这块感觉没什么好讲的都是流程，需要看到的就是，使用了两次核函数完成topk任务，正如我们所想，第一步输出中间结果，中间结果是每个block块中的前topk个，输出大小为（topk*GRID_SIZE），第二部将中间结果当作输入求最终输出结果，最终结果是在1个块中的前topk个。

那么来看主要的核函数，以及设备函数。

__device__ __host__ void insert_sort(int*array,int k,int data) {
	//由__device__ __host__修饰符用于一个函数，表示该函数可以在GPU（设备）和CPU（主机）上执行
	//__device__ 声明一个函数为设备函数，该函数只能在GPU上执行，一般是被其他设备函数或者kernel调用
	// __host__ 声明一个函数为主机函数，该函数只能在CPU上执行，通常用于管理设备内存、启动核心等操作。
	for (int i = 0; i < k; i++) {
		//如果数据重复，就不参与排序了，选的是前20个最大的。
		if (array[i] == data) {
			return;
		}
	}
	//如果要插入的元素小于数组最后一个元素，那就不参与排序，因为就不是前20个最大的了
	if (data < array[k - 1]) {
		return;
	}
	//从倒数第二个开始，向前比较，如果当前数据比数据元素大，那么数组元素就向后移位，如果小于则将该位置的后一位作为插入位置。位置
	for (int i = k - 2; i >= 0; i--) {
		if (data > array[i]) {
			array[i + 1] = array[i];
		}
		else {
			array[i + 1] = data;
			return;
		}
	}

	//如果data比所有数据都大，那么将这个数据插入到改数组的第一个位置
	array[0] = data;
}

首先第一个函数insert_sort函数,实现一个插入排序，但其是由__device__,__host__两个关键字共同限制，表示既为设备函数，也为主机函数，在主机端和设备端都可调用，通常用于设备端与主机端的比较。

cuda中最常用的关键字，__device__，__host__，和 __global__ 是三个用于指定函数类型和执行空间的关键字。其中device定义的函数，一般被global定义的核函数或者其他device定义的设备函数调用。而Host定义的函数只能在CPU上也就是主机端code被调用，而不是由__device__,或者—__global__限定的函数。__global__ 关键字用于声明CUDA内核函数，这是一种特殊的函数，可以由CPU调用并在GPU上并行执行。这3个关键字中只有device与host可以组队，而global则不行，有点像明明是三个人的电影，我却始终不能有姓名。

然后这个函数吗原理也没啥好讲的，应该可以看得懂，注释的很详细，排除重复，以及小于数组最后一个数，剩下就开始将数据与数组元素比较，插入到合适的位置。

接下来就是重点核函数：

	//申请共享内存数据，用与保存每个块的计算结果
	__shared__ int ken[BLOCK_SIZE * topk];

首先可以看到核函数一进来就申请了一个共享内存数组。使用的__shared__关键在在kernel内声明。首先需要清楚的是，共享内存是在每个block中都有的，这里定义的的共享内存数组ken，在每个线程块中都有，块内所有线程均可访问，即共享内存，用于存储该块的top-k结果。其中BLOCK_SIZE是每个线程块的线程数。

	int top_array[topk];
	//top_array初始化，给最小值
	for (int i = 0; i < topk; i++) {
		top_array[i] = INT_MIN;
	}

申请了一个局部的top_arry数组，用于存储当前线程处理的top-k结果，并赋了初值。这是每个线程私有的，一般位于寄存器，如果溢出就位于局部内存（即local）。

	//插入排序
	//对数组中的所有数据进行插入排序
	for (int idx = threadIdx.x + blockDim.x * blockIdx.x; idx < length; idx += gridDim.x * blockDim.x) {
		insert_sort(top_array, topk, input[idx]);
	}

接着使用了一个for循环，这个for循环的初值是线程的全局索引，条件是小于数组长度，增量为网格（grid）中的所有线程数。这意味着每个线程在处理完其当前元素后，会跳过整个网格中的所有其他线程所处理的元素，直接处理下一个它需要处理的元素。

通过这种方式，该循环确保如果数组 input 的长度大于网格中线程的总数，每个线程将能够处理数组中的多个元素。每个线程处理的元素索引之间的间隔等于整个网格的线程总数，从而实现了数据的均匀分配。(这种方法称作网格跨步循环，用于当需要操作的元素数大于核函数的总线程数时。)

这里可能有人会疑问，我这不是就只有 gridDim.x * blockDim.x个线程么，怎么全局索引还要加这么多线程，这里就需要提到两个重要概念：

线程索引：记住线程索引不是线程本身。线程索引是一个数值，它唯一地标识了在CUDA网格中的每个线程。这个索引是根据线程在网格、线程块和线程内的位置计算出来的。线程索引用于确定每个线程应该处理的数据元素。在处理大型数据集时，这使得每个线程可以独立地工作在数组或内存中的不同位置。但是在编程时，其实线程索引与线程本身很容易混淆，二者这么看来是有重叠，但又不一样。

线程身份：线程身份是指线程的物理或逻辑存在。在CUDA中，线程通过它们的 threadIdx 和 blockIdx 被识别，但这些标识符并不提供关于线程在GPU硬件上的具体物理位置或身份的信息。在CUDA程序设计中，通常不需要知道线程的具体物理身份。相反，重要的是如何利用线程索引来合理地分配和控制数据处理。

在CUDA中，线程索引是核心概念，因为它决定了线程如何访问和处理数据。线程身份（如物理ID）在日常的CUDA编程中不是一个重点，因为CUDA的设计允许开发者从更高的抽象层面上考虑并行计算问题，而无需关注底层硬件的具体细节。

所以需要理解的是，当在调用核函数时，需要核函数的执行配置，GRID_SIZE,BLOCK_SIZE,这两个代表的这抽象的网络和块，当然也包括线程，抽象的是什么意思，即就是物理逻辑不存在，硬件上不是按着这么划分的。在核函数内，核函数执行的线程数是一定的，但至于是我前一个循环的线程是不是我当前循环所使用的线程，那没有意义，我只要确保我核函数内每回执行时线程数一致即可。

这里在用数据去感受一下，首先数组大小是N = 100000000，BLOCK_SIZE,GRID_SIZE是256，32。

总线程数为8192，那么需要多少次for循环呢，12208次。所以要使用8192个线程处理N个数组，需要12208次循环。

在来深入理解以下，比如第一个循环线程索引从0~8191，下一循环线程索引就从8192-16384，一次类推，所以线程索引只是线程索引，并不代表线程数的多少，在核函数内并行执行的线程数是一致的。

	//维护好的top array放进共享内存数组
	for (int i = 0; i < topk; i++) {
		ken[topk * threadIdx.x + i] = top_array[i];
	}
	__syncthreads();

然后将每个线程处理后的数组的top array放进共享内存数组ken中。ken的大小是块内线程总数*topk。每个线程单独维护着top array，也就是每个线程都有一个数组top_array。这也能解释在赋值时ken的索引为topk*threaIdx.x+i，因为threaIdx.x是块内线程编号，每个线程都有大小为topk的数组top_array,将他们顺序放入即为topk*threadx.x+i,这里是做一个线程块内多线程数据合并。

	//共像内存中的数据合并，并行归约。
	// 每一步都将当前活动的线程数减半，这些线程合并相邻的 top_array。
	// 这个过程在每个线程块内部进行，最终得到该块的局部前 k 个最大值。
	for (int i = BLOCK_SIZE/2; i >= 1; i /= 2) {
		if (threadIdx.x < i) {
			for (int m = 0; m < topk; m++) {
				insert_sort(top_array, topk, ken[topk * (threadIdx.x + i) + m]);
			}
		}
		__syncthreads();
		if (threadIdx.x < i) {
			for (int m = 0; m < topk; m++) {				
				ken[topk * threadIdx.x + m] = top_array[m];
			}
		}
		__syncthreads();
	}

接下来的操作是一个归约操作，每一步线程数减半，并且合并块内共享内存数组ken中的元素顺序。这种归约模式是高效并行算法的一个典型例子，它利用了线程之间的协作，将多个线程的结果合并成一个更小的集合。在每一步中，活跃的线程数量减半，最终得到每个线程块的局部Top-K结果。

这里需要重点记住的是，Ken是整个块内的线程所共享的数组，所以对Ken的索引需要使用线程索引，而top_array是每个线程私有的，所以不需要加线程索引，在核函数内自己就执行了。

在来深入分析以下，这是怎么做到的。

ken[topk * (threadIdx.x + i) + m]

首先线程索引，在大循环中一次减半，也就是说在这个for循环中用的线程数是依次减少的。

ken数组的大小5120。

i的变化是128，64，32，16，8，4，2，1。

threadIdx.x的变化，从0~127，到0~63，到0~31，到0~15，到0~7，到0~3，到0~1，最后0。

而threadIdx.x是与i共同变化的。

m的变化是0~19。

试着去算算首先是大循环下第一次循环数组索引范围是多少。最小值：20*（0+128）+0 = 2560，最大值：20*（127+128）+19 = 5119，表示对ken数组的后一半数据，将这些数组插入哪里呢，插入到前128个线程（即0~127）维护的top_array数组中，这里表示插入到前128个线程所维护的top_array数组。

然后后面再将前128个线程所维护的top_array分别在原位置更新Ken。其余依次类推。最终ken中的前20个就保留了block块内最大的topk数组。而由于ken是一个block私有的，因此如果有多个块block，最后output就包含多个块的ken。最终只取每个ken的前20个值即可。

最后：将ken

	//将最终结果写入输出数组，只使用每个线程块一个线程，可以是0，也可以使其他，
	// 用于将前topk个最大值写入输出数组相应的位置
	if (blockIdx.x * blockDim.x < length) {
		if (threadIdx.x == 0) {
			for (int i = 0; i < topk; i++) {
				output[topk * blockIdx.x + i] = ken[i];
			}
		}
	}

此外还有cpu的topk函数

void cpu_topk(int* input, int* output, int length, int k) {
	for (int i = 0; i < length; i++) {
		insert_sort(output, k, input[i]);
	}
}

这也没啥讲的。

然后总结一下：

• 核函数内声明不带关键字的变量或数组是每个线程私有的。
• 对于线程数小于要操作的数据可以采用网格跨步循环。附上一个小代码。

__global__ void MyKernel(int *array, int arraySize) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;

    for (int i = idx; i < arraySize; i += stride) {
        // 处理array[i]
    }
}

• 要注意不同类型数据的访问，全局内存的数据，一般采用全局索引，共享内存的索引一般采用块内索引。
• 使用规约算法处理大数据时非常有用，类似于递归，可以灵活地调整线程块的大小和数量，而不受处理的数据量的限制。它确保了即使数组大小远远超过线程总数，每个元素也都能被相应的线程处理。
• 理解线程索引与线程身份的含义不同，在核函数内执行的线程索引并不一定代表线程本身，线程，block，grid，在cuda编程中是一种抽象的上层逻辑，而不是底层硬件的物理逻辑。底层是通过流式处理器SM，以及许多cuda core，实现的，包含处理int，float，double等数据类型的加减乘除运算。

最终附上完全代码。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include  "device_functions.h"
#include <stdio.h> //c标准输出输出库
#include <math.h>
#include <stdlib.h> //c标准库包含rand，产生随机数

//topk 问题 数组前k个大的元素
//归约，累加求和

#define N 100000000  //数据大小
#define BLOCK_SIZE 256  //一个块中有256个线程
#define GRID_SIZE 32  //32 每个网格中有32个块
#define topk 20 

__managed__ int source[N];   //原数组 
//__managed__  cuda关键字，用于声明所谓的托管内存，允许内存在CPU和GPU之间自动共享。
//用 __managed__ 声明的变量可以同时被 CPU 和 GPU 访问，无需手动在主机（CPU）和设备（GPU）之间复制数据。
//使用托管内存简化了内存管理，因为它允许 CPU 和 GPU 在无需显式数据传输命令的情况下访问相同的内存。

__managed__ int gpu_result[topk];  //topk最终结果
__managed__ int _1_pass_result[topk * GRID_SIZE];//每个block的前20个，即中间结果

//理论，求一个大数组的前20个最大值，先将数组放入GPU内，每个block中求出最大的前20个值，放入_1_passresult
//然后每个block前20个值放一块在求前20个值得到最中结果


__device__ __host__ void insert_sort(int*array,int k,int data) {
	//由__device__ __host__修饰符用于一个函数，表示该函数可以在GPU（设备）和CPU（主机）上执行
	//__device__ 声明一个函数为设备函数，该函数只能在GPU上执行，一般是被其他设备函数或者kernel调用
	// __host__ 声明一个函数为主机函数，该函数只能在CPU上执行，通常用于管理设备内存、启动核心等操作。
	for (int i = 0; i < k; i++) {
		//如果数据重复，就不参与排序了，选的是前20个最大的。
		if (array[i] == data) {
			return;
		}
	}
	//如果要插入的元素小于数组最后一个元素，那就不参与排序，因为就不是前20个最大的了
	if (data < array[k - 1]) {
		return;
	}
	//从倒数第二个开始，向前比较，如果当前数据比数据元素大，那么数组元素就向后移位，如果小于则将该位置的后一位作为插入位置。位置
	for (int i = k - 2; i >= 0; i--) {
		if (data > array[i]) {
			array[i + 1] = array[i];
		}
		else {
			array[i + 1] = data;
			return;
		}
	}

	//如果data比所有数据都大，那么将这个数据插入到改数组的第一个位置
	array[0] = data;
}

__global__ void gpu_topk(int* input, int* output, int length, int k) {
	/*
	参数说明：
	input 输入数组
	output 输出数组
	length 输入数组的长度
	*/
	//申请共享内存数据，用与保存每个块的计算结果
	__shared__ int ken[BLOCK_SIZE * topk];
	

	int top_array[topk];
	//top_array初始化，给最小值
	for (int i = 0; i < topk; i++) {
		top_array[i] = INT_MIN;
	}

	//插入排序
	//对数组中的所有数据进行插入排序
	for (int idx = threadIdx.x + blockDim.x * blockIdx.x; idx < length; idx += gridDim.x * blockDim.x) {
		insert_sort(top_array, topk, input[idx]);
	}

	//维护好的top array放进共享内存数组
	for (int i = 0; i < topk; i++) {
		ken[topk * threadIdx.x + i] = top_array[i];
	}
	__syncthreads();

	//共像内存中的数据合并，并行归约。
	// 每一步都将当前活动的线程数减半，这些线程合并相邻的 top_array。
	// 这个过程在每个线程块内部进行，最终得到该块的局部前 k 个最大值。
	for (int i = BLOCK_SIZE/2; i >= 1; i /= 2) {
		if (threadIdx.x < i) {
			for (int m = 0; m < topk; m++) {
				insert_sort(top_array, topk, ken[topk * (threadIdx.x + i) + m]);
			}
		}
		__syncthreads();
		if (threadIdx.x < i) {
			for (int m = 0; m < topk; m++) {				
				ken[topk * threadIdx.x + m] = top_array[m];
			}
		}
		__syncthreads();
	}

	//将最终结果写入输出数组，只使用每个线程块一个线程，可以是0，也可以使其他，
	// 用于将前topk个最大值写入输出数组相应的位置
	if (blockIdx.x * blockDim.x < length) {
		if (threadIdx.x == 0) {
			for (int i = 0; i < topk; i++) {
				output[topk * blockIdx.x + i] = ken[i];
			}
		}
	}

}

void cpu_topk(int* input, int* output, int length, int k) {
	for (int i = 0; i < length; i++) {
		insert_sort(output, k, input[i]);
	}
}

int main(){

	//为原数组赋初值
	printf("初始化源数据.....\n");
	for (int i = 0; i < N; i++) {
		source[i] = rand();
	}
	printf("完成初始化源数据.....\n");

	//cuda事件-计时
	cudaEvent_t start, stop_gpu, stop_cpu;
	cudaEventCreate(&start);
	cudaEventCreate(&stop_gpu);
	cudaEventCreate(&stop_cpu);
	cudaEventRecord(start);
	cudaEventSynchronize(start);//事件同步
	//这个函数用于等待一个 CUDA 事件完成。
	// 当你在 CUDA 程序中设置一个事件时，比如 cudaEventRecord(event, stream)，
	// 它会在特定的流（stream）中标记一个点。cudaEventSynchronize(event) 会阻塞调用线程，
	// 直到该事件发生，即直到 GPU 上的相关操作完成。
	printf("GPU Run *************\n");
	int times = 1;
	//计算
	for (int i = 0; i < times; i++) {
		gpu_topk << <GRID_SIZE, BLOCK_SIZE >> > (source, _1_pass_result, N, topk);
		gpu_topk << <1, BLOCK_SIZE >> > (_1_pass_result, gpu_result, topk * GRID_SIZE, topk);
		cudaDeviceSynchronize();
		//cudaDeviceSynchronize() 函数会阻塞调用线程，直到 GPU 完成所有队列中的操作。
		// 这包括所有 CUDA 核心、内存复制和其他相关的 GPU 操作。
	}
	printf("GPU Run Complete %d 次*************\n",times);
	cudaEventRecord(stop_gpu);
	cudaEventSynchronize(stop_gpu);


	//cpu结果初始化
	int cpu_result[topk] = { 0 }; //cpu结果存储
	printf("CPU Run *************\n");
	//计算
	cpu_topk(source, cpu_result, N, topk);
	printf("GPU Run Complete *************\n");
	cudaEventRecord(stop_cpu);
	cudaEventSynchronize(stop_cpu);

	//计算两次时间
	float time_cpu, time_gpu;
	cudaEventElapsedTime(&time_gpu, start, stop_gpu);
	cudaEventElapsedTime(&time_cpu, stop_gpu, stop_cpu);

	//判断GPU计算是否有误
	bool error = false;
	for (int i = 0; i < topk; i++) {
		printf(" CPU top%d: %d; Gputop%d: %d;\n", i + 1, cpu_result[i], i + 1, gpu_result[i]);
		if (fabs(gpu_result[i] - cpu_result[i]) > 0) {
			error = true;
		}
	}
	printf("Result:%s\n", (error ? "Error" : "pass"));
	printf("CPU time: %.2f; GPU time: %.2f\n", time_cpu, time_gpu);

	return 0;
}