北京白癜风哪治的好 http://pf.39.net/bdfyy/bdfzj/在前三部分中我们介绍了CUDA开发的大部分基础知识，例如启动内核来执行并行任务、利用共享内存来执行快速归并、将可重用逻辑封装为设备函数以及如何使用事件和流来组织和控制内核执行。
本文是本系列的最后一部分，我们将讨论原子指令，它将允许我们从多个线程中安全地操作同一内存。我们还将学习如何利用这些操作来创建互斥锁，互斥锁是一种编码模式，它允许我们“锁定”某个资源，以便每次只由一个线程使用它。可以说这两个概念是任何多线程的基础。
还是从头开始，我们导入和加载库，并确保有一个GPU。
importwarningsfromdatetimeimportdatetimefromtimeimportperf_counterimportmatplotlibasmplimportmatplotlib.pyplotaspltimportmatplotlib.tickerastickerimportnumpyasnpimportrequestsimportnumbafromnumbaimportcudafromnumba.core.errorsimportNumbaPerformanceWarningfromtqdm.autoimporttrangeprint(np.__version__)print(numba.__version__)print(mpl.__version__)#IgnoreNumbaPerformanceWarningwarnings.simplefilter("ignore",category=NumbaPerformanceWarning)#1.21.6#0.56.2#.5.#Found1CUDAdevices#id0bTeslaT4[SUPPORTED]#ComputeCapability:7.5#PCIDeviceID:4#PCIBusID:0#UUID:GPU-5ffb66-a0f2-a5b0-15a-18dd9ba#Watchdog
isabled#FP2/FP64PerformanceRatio:2#Summary
1/1devicesaresupported#True
原子操作
GPU编程的思想是基于尽可能多地并行执行相同的指令。对于许多可以并行任务，线程之间不需要合作或使用其他线程使用的资源，只要保证自己运行正确即可。但是一些需要同步执行的操作，例如归并（Reduce），需要通过算法的设计确保相同的资源只被线程的子集使用，所以需要通过使用同步线程确保所有其他线程都是最新的。
在某些情况下，多个线程必须对同一个数组进行读写。当试图同时执行读或写操作时，这可能会导致问题，例如假设我们有一个将一个值加1的内核。
#Example4.1:Adataracecondition.
cuda.jitdefadd_one(x)
[0]=x[0]+1
当我们用一个线程块启动这个内核时，我们将在输入数组中存储一个值1。
dev_val=cuda.to_device(np.zeros((1,)))add_one[1,1](dev_val)dev_val.copy_to_host()#array([1.])
如果我们启动10个区块，每个区块有16个线程时会发生什么?10×16×1加到同一个内存元素中，所以我们应该希望dev_val中得到的值是。对吧?
dev_val=cuda.to_device(np.zeros((1,)))add_one[10,16](dev_val)dev_val.copy_to_host()
如果使用上面的操作，我们不太可能在dev_val中得到到。为什么呢?因为线程同时在读写同一个内存变量!
下面是当四个线程试图从同一个全局内存中读写时可能发生的情况的示意图。线程1-从全局寄存器读取相同的值0的次数不同(t分别为0,2,2)。它们都增加1，并在t=4,7和8时写回全局内存。线程4开始的时间比其他线程稍晚，在t=5时。此时，线程1已经写入全局内存，因此线程4读取的值为1。它最终会在t=12时将全局变量改写为2。
从同一个全局内存中读写的多个线程的情况示意图，也就是说这个操作是非线程安全的
当一个线程对内容进行操作时，资源被禁止读/写，所以确保每个线程在读时获得更新的值，而其他线程看到它的写。这种线程安全的操作通常较慢。
如果我们想要获得最初期望的结果(如图2所示)，我们应该用原子操作（上面的线程安全的操作）替换非原子加法操作。原子操作将确保无论读/写的是什么内存，每次都由单个线程完成。
#Example4.2:Anatomicaddwithoutraceconditions.
cuda.jitdefadd_one_atomic(x):cuda.atomic.add(x,0,1)#Argumentsarearray,arrayindex,valuetoadddev_val=cuda.to_device(np.zeros((1,)))add_one_atomic[10,16](dev_val)dev_val.copy_to_host()#array([.])
原子加法操作示例:计算直方图
为了更好地理解在哪里以及如何使用原子操作，我们将使用直方图计算。假设有人想数一数在某一文本中字母表中的每个字母有多少个。实现这一目标的一个简单算法是创建26个“桶”，每个桶对应英语字母表中的一个字母。然后我们将遍历文本中的字母，每当我们遇到“a”时，我们将增加第一个bucket1，每当我们遇到“b”时，我们将增加第二个bucket1，以此类推。
在标准Python中，可以使用字典来实现我们的“桶”，每个字典都将一个字母与一个数字联系起来。由于我们是在GPU上进行操作，所以这里将使用数组代替字典，并且将存储所有个ASCII字符，而不是存储26个字母。
在此之前，我们需要将字符串转换为“数字”数组。在这种情况下可以将UTF-8字符串转换为uint8数据类型。
defstr_to_array(x):returnnp.frombuffer(bytes(x,"utf-8"),dtype=np.uint8)defgrab_uppercase(x):returnx[65:65+26]defgrab_lowercase(x):returnx[97:97+26]my_str="CUDAbyNumbaExamples"my_str_array=str_to_array(my_str)#array([67,85,68,65,2,98,,2,78,,,98,97,#2,69,,97,,,,,],dtype=uint8)
小写字母和大写字母有不同的代码。我们将用函数来实现仅选择小写字母或仅选择大写字母。
并且Numpy已经提供了一个直方图函数，我们将使用它来验证结果并比较运行时。
histo_np,bin_edges=np.histogram(my_str_array,bins=,range=(0,))np.testing.assert_allclose(bin_edges,np.arange())#Binedgesare1morethanbinsdefplot_letter_histogram(hist,bin_edges,kind="percent",ax=None):width=0.8start=bin_edges[0]stop=bin_edges[0]+hist.shape[0]ifaxisNone:ax=plt.gca()ax.bar(np.arange(start,stop),hist,width=width)ax.xaxis.set_major_locator(ticker.MultipleLocator(1))ax.xaxis.set_major_formatter(ticker.FuncFormatter(lambdax,pos
"{int(x):c}"))ax.set(xlim=[start-width,stop-1+width],ylabel=kind.title())ifkind=="count":ax.yaxis.set_major_formatter(ticker.FuncFormatter(lambdax,pos"{x:.0f}"))else:sum_hist=hist.sum()ifkind=="probability":ax.yaxis.set_major_formatter(ticker.FuncFormatter(lambdax,pos"{x/sum_hist:.2f}"))else:ax.yaxis.set_major_formatter(ticker.FuncFormatter(lambdax,pos"{x/sum_hist:.0%}"))fig,axs=plt.subplots(1,2,figsize=(10,),sharey=True)plot_letter_histogram(grab_lowercase(histo_np),grab_lowercase(bin_edges),kind="count",ax=axs[0])plot_letter_histogram(grab_uppercase(histo_np),grab_uppercase(bin_edges),kind="count",ax=axs[1])axs[0].set(title="Lowercase")axs[1].set(title="Uppercase");
首先让我们编写CPU版本的函数来理解机制。
defhistogram_cpu(arr):histo=np.zeros(,dtype=np.int64)forcharinarr:ifchar:histo[char]+=1returnhistohisto_cpu=histogram_cpu(my_str_array)assert(histo_cpu-histo_np).sum()==0#Matchesnumpyversion
由于每个ASCII字符都映射到元素数组中的一个bin，因此我们需要做的就是找到它的bin并加1，只要该bin在0和（包括）之间即可。
理解了我们的函数实现，下面我们来完成GPU版本：
#Example4.:AGPUhistogram
cuda.jitdefkernel_histogram(arr,histo):i=cuda.grid(1)threads_per_grid=cuda.gridsize(1)foriarrinrange(i,arr.size,threads_per_grid):ifarr[iarr]:cuda.atomic.add(histo,arr[iarr],1)cuda.jitdefkernel_zero_init(arr):i=cuda.grid(1)threads_per_grid=cuda.gridsize(1)foriarrinrange(i,arr.size,threads_per_grid):arr[iarr]=0threads_per_block=blocks_per_grid=2my_str_array_gpu=cuda.to_device(my_str_array)histo_gpu=cuda.device_array((,),dtype=np.int64)kernel_zero_init[1,](histo_gpu)kernel_histogram[blocks_per_grid,threads_per_block](my_str_array_gpu,histo_gpu)histo_cuda=histo_gpu.copy_to_host()assert(histo_cuda-histo_cpu).sum()==0
可以看到我们的函数可以正常工作了。这个内核非常简单并且与串行版本结构相同。它以标准的1D循环结构开始，使用原子加法。Numba中的原子加法有三个参数：需要递增的数组(histo)、需要加法操作的数组位置（arr[iarr]），需要相加的值（在本例中为1）。
下面让我们看看在更大的数据集中如何使用
#Getthe
