从头开始进行CUDA编程:Numba并行编程的基本概念

PU(图形处理单元)最初是为计算机图形开发的,但是现在它们几乎在所有需要高计算吞吐量的领域无处不在。这一发展是由GPGPU(通用GPU)接口的开发实现的,它允许我们使用GPU进行通用计算编程。这些接口中最常见的是CUDA,其次是OpenCL和最近刚出现的HIP。

Python中使用CUDA

CUDA最初被设计为与C兼容后来的版本将其扩展到c++和Fortran。在Python中使用CUDA的一种方法是通过Numba,这是一种针对Python的即时(JIT)编译器,可以针对gpu(它也针对cpu,但这不在我们讨论的范围内)。Numba为我们提供了一个可以直接使用Python子集,Numba将动态编译Python代码并运行它。虽然它没有实现完整的CUDA API,但与cpu相比它支持的特性已经可以帮助我们进行并行计算的加速。

Numba并不是唯一的选择。CuPy 提供了通过基于CUDA的并且兼容Numpy的高级函数,PyCUDA提供了对CUDA API更细粒度的控制,英伟达也发布了官方CUDA Python。

本文不是 CUDA 或 Numba 的综合指南,本文的目标是通过用Numba和CUDA编写一些简单的示例,这样可以让你了解更多GPU相关的知识,无论是是不是使用Python,甚至C编写代码,它都是一个很好的入门资源。

GPU 的并行编程简介

GPU 相对于 CPU 的最大优势是它们能够并行执行相同的指令。单个 CPU 内核将一个接一个地串行运行指令。在 CPU 上进行并行化需要同时使用其多个内核(物理或虚拟)。例如一般的计算机有 4-8 个内核,而GPU 拥有数千个计算核心。有关这两者的比较,请参见下面的图 1。GPU 内核通常速度较慢,且只能执行简单的指令,但它们的数量通常可以弥补这些缺点。

GPU 编程有四个主要方面问题:

1、理解如何思考和设计并行的算法。因为一些算法是串行设计的,把这些算法并行化可能是很困难的。

2、学习如何将CPU上的结构(例如向量和图像)映射到 GPU 上例如线程和块。循环模式和辅助函数可以帮助我们解决这个问题。

3、理解驱动 GPU 编程的异步执行模型。不仅 GPU 和 CPU 相互独立地执行指令,GPU的流还允许多个处理流在同一个GPU上运行,这种异步性在设计最佳处理流时非常重要。

4、抽象概念和具体代码之间的关系:这是通过学习 API 及其细微差别来实现的。

上图为简化的CPU架构(左)和GPU架构(右)。计算发生在ALU(算术逻辑单元)中,DRAM保存数据,缓存保存的数据可以更快地访问,但通常容量更小。

开始编写代码

这里的环境要求是:Numba版本> 0.55和一个GPU。

 import numpy as np
 import numba
 from numba import cuda
 
 print(np.__version__)
 print(numba.__version__)
 
 cuda.detect()
 
 # 1.21.6
 # 0.55.2
 #
 # Found 1 CUDA devices
 # id 0             b'Tesla T4'                              [SUPPORTED]
 #                       Compute Capability: 7.5
 #                            PCI Device ID: 4
 #                               PCI Bus ID: 0
 #                                     UUID: GPU-e0b8547a-62e9-2ea2-44f6-9cd43bf7472d
 #                                 Watchdog: Disabled
 #              FP32/FP64 Performance Ratio: 32
 # Summary:
 # 1/1 devices are supported

Numba CUDA的主要操作时是CUDA.jit的装饰器,它定义函数将在GPU中运行。

我们首先写一个简单的函数,它接受两个数字相加然后将它们存储在第三个参数的第一个元素上。

 # Example 1.1: Add scalars
 @cuda.jit
 def add_scalars(a, b, c):
     c[0] = a + b
 
 dev_c = cuda.device_array((1,), np.float32)
 
 add_scalars[1, 1](2.0, 7.0, dev_c)
 
 c = dev_c.copy_to_host()
 print(f"2.0 + 7.0 = {c[0]}")
 #  2.0 + 7.0 = 9.0

因为GPU只能处理简单的工作,所以我们的内核只运行一个函数,后面会将这个函数称之为内核。

第一个需要注意的是内核(启动线程的GPU函数)不能返回值。所以需要通过传递输入和输出来解决这个问题。这是C中常见的模式,但在Python中并不常见。

在调用内核之前,需要首先在设备上创建一个数组。如果想要显示返回值则需要将它复制回CPU。这里就有一个隐形的问题:为什么选择float32(单精度浮点数)?这是因为虽然大多数GPU都支持双精度运算,但双精度运算的时间可能是单精度运算的4倍甚至更长。所以最好习惯使用np.float32和np.complex64而不是float / np.float64和complex / np.complex128

我们的函数定义与普通的函数定定义相同,但调用却略有不同。它在参数之前有方括号:add_scalars[1, 1](2.0, 7.0, dev_c)

这些方括号分别表示网格中的块数和块中的线程数,下面使用CUDA进行并行化时,会进一步讨论。

使用CUDA进行并行化编程

CUDA网格

当内核启动时它会得到一个与之关联的网格,网格由块组成;块由线程组成。下图2显示了一维CUDA网格。图中的网格有4个块。网格中的块数保存在一个特殊的变量中,该变量可以在内核中通过gridDim.x直接访问,这里x是指网格的第一维度(在本例中是唯一的维度)。二维网格也有通过y还有三维网格z变量来访问。到目前2022年,还没有四维网格或更高的网格。在内核内部可以通过使用 blockIdx.x 找出正在执行的块,例如我们这个例子它将从 0 运行到 3。

每个块都有一定数量的线程,保存在变量blockDim.x中。线程索引保存在变量 threadIdx.x 中,在这个示例中变量将从 0 运行到 7。

不同块中的线程被安排以不同的方式运行,访问不同的内存区域并在其他一些方面有所不同,本文主要介绍简单的入门所以我们将跳过这些细节。

当我们在第一个示例中使用参数[1,1]启动内核时,我们告诉CUDA用一个线程运行一个块。通过修改这两个值可以使用多个块和多现线程多次运行内核。 threadIdx.x 和 blockIdx.x 每个线程的唯一标识。

下面我们对两个数组求和,这比对两个数字求和复杂:假设每个数组都有20个元素。如上图所示,我们可以用每个块8个线程启动内核。如果我们希望每个线程只处理一个数组元素,那么我们至少需要4个块。启动4个块,每个块8个线程,我们的网格将启动32个线程。

对于多线程处理,最需要弄清楚是如何将线程下标映射到数组下标(因为每个线程要独立处理部分数据)。threadIdx.x 从 0 运行到 7,因此它们自己无法索引我们的数组,不同的块也具有相同的threadIdx.x。但是他们有不同的blockIdx.x。为了获得每个线程的唯一索引,我们需要组合这些变量:

 i = threadIdx.x + blockDim.x * blockIdx.x

对于第一个块,blockIdx.x = 0,i 将从 0 运行到 7。对于第二个块,blockIdx.x = 1。由于 blockDim.x = 8,将得到 8 运行到 15。类似地,对于 blockIdx. x = 2,将从 16 跑到 23。在第四个也是最后一个区块中,将从 24 跑到 31。见下表 1。

这样虽然将每个线程映射到数组中的每个元素……但是现在我们遇到了一些线程会溢出数组的问题,因为数组有 20 个元素,而 i 的最大值是 32-1。解决方案很简单:对于那些溢出线程,不要做任何事情!

 # Example 1.2: Add arrays
 @cuda.jit
 def add_array(a, b, c):
     i = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
     if i < a.size:
         c[i] = a[i] + b[i]
 
 N = 20
 a = np.arange(N, dtype=np.float32)
 b = np.arange(N, dtype=np.float32)
 dev_c = cuda.device_array_like(a)
 
 add_array[4, 8](a, b, dev_c)
 
 c = dev_c.copy_to_host()
 print(c)
 #  [ 0.  2.  4.  6.  8. 10. 12. 14. 16. 18. 20. 22. 24. 26. 28. 30. 32. 34. 36. 38.]

在较新版本的 Numba 中可能会会收到一条警告,指出我们使用内核使用了非设备上的数据。这条警告的产生的原因是将数据从主机移动到设备非常慢, 我们应该在所有参数中使用设备数组调用内核。所以我们需要预先将数组从主机移动到设备:

 dev_a = cuda.to_device(a)
 dev_b = cuda.to_device(b)

每个线程唯一索引的计算可能很快就会过期, Numba 提供了非常简单的包装器 cuda.grid,它以网格维度作为唯一参数调用。所以我们新的内核将如下:

 # Example 1.3: Add arrays with cuda.grid
 @cuda.jit
 def add_array(a, b, c):
     i = cuda.grid(1)
     if i < a.size:
         c[i] = a[i] + b[i]
 
 add_array[4, 8](dev_a, dev_b, dev_c)
 
 c = dev_c.copy_to_host()
 print(c)
 
 #  [ 0.  2.  4.  6.  8. 10. 12. 14. 16. 18. 20. 22. 24. 26. 28. 30. 32. 34. 36. 38.]

如果我们改变数组的大小时会发生什么?我们这里不改变函数而更改网格参数(块数和每个块的线程数),这样就相当于启动至少与数组中的元素一样多的线程。

设置这些参数有一些”科学“和一些”艺术“。对于“科学”:(a)它们应该是 2 的倍数,通常在 32 到 1024 之间,并且(b)它们应该都被使用以最大化占用率(有多少线程同时处于活动状态 )。所以Nvidia 提供了一个可以帮助计算这些的电子表格。(https://docs.nvidia.com/cuda/cuda-occupancy-calculator/index.html) 对于“艺术”而言,没有什么可以预测内核的行为,因此如果真的想优化这些参数,需要根据不同的输入来分析代码。但是根据经验一般GPU 的“合理”线程数是 256。

 N = 1_000_000
 a = np.arange(N, dtype=np.float32)
 b = np.arange(N, dtype=np.float32)
 
 dev_a = cuda.to_device(a)
 dev_b = cuda.to_device(b)
 dev_c = cuda.device_array_like(a)
 
 threads_per_block = 256
 blocks_per_grid = (N + (threads_per_block - 1)) // threads_per_block
 # Note that
 #     blocks_per_grid == ceil(N / threads_per_block)
 # ensures that blocks_per_grid * threads_per_block >= N
 
 add_array[blocks_per_grid, threads_per_block](dev_a, dev_b, dev_c)
 
 c = dev_c.copy_to_host()
 np.allclose(a + b, c)
 
 #  True

代码方面是没有问题了,但是因为硬件限制GPU 不能运行任意数量的线程和块。通常每个块不能超过 1024 个线程,一个网格不能超过 2¹⁶ - 1 = 65535 个块。但是这并不是说可以启动 1024 × 65535 个线程……因为还需要根据寄存器占用的内存量以及其他的因素考虑。还有一点就是处理不适合 GPU RAM 的大型数组(也就是OOM)。在这些情况下,可以使用多个 GPU 分段处理数组(单机多卡)。

在 Python 中,硬件限制可以通过 Nvidia 的 cuda-python 库的函数 cuDeviceGetAttribute 获得,具体请查看该函数说明。

Grid-stride循环

在每个网格的块数超过硬件限制但显存中可以容纳完整数组的情况下,可以使用一个线程来处理数组中的多个元素,这种方法被称为Grid-stride。除了克服硬件限制之外,Grid-stride还受益于重用线程,这样可以将线程创建/销毁开销降到最低。Mark Harris 的文章 CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops 详细介绍这个方法,这里就不再介绍了。

在 CUDA 内核中添加一个循环来处理多个输入元素,这个循环的步幅等于网格中的线程数。这样如果网格中的线程总数 (threads_per_grid = blockDim.x * gridDim.x) 小于数组的元素数,则内核处理完索引 cuda.grid(1)它将处理索引 cuda.grid(1) + threads_per_grid,直到处理完所有数组元素,我们来看代码。

 # Example 1.4: Add arrays with grid striding
 @cuda.jit
 def add_array_gs(a, b, c):
     i_start = cuda.grid(1)
     threads_per_grid = cuda.blockDim.x * cuda.gridDim.x
     for i in range(i_start, a.size, threads_per_grid):
         c[i] = a[i] + b[i]
 
 threads_per_block = 256
 blocks_per_grid_gs = 32 * 80  # Use 32 * multiple of streaming multiprocessors
 # 32 * 80 * 256 < 1_000_000 so one thread will process more than one array element
 
 add_array_gs[blocks_per_grid_gs, threads_per_block](dev_a, dev_b, dev_c)
 c = dev_c.copy_to_host()
 np.allclose(a + b, c)
 
 #  True

这段代码与上面的代码非常相似,不同之处在于从cuda.grid(1)开始,但是执行更多的元素,每个threads_per_grid执行一个元素,直到到达数组的末尾。

那么那种方式更快呢?

CUDA 内核的计算时间

GPU 编程的目标就是提高速度。因此准确测量代码执行时间非常重要。

CUDA内核是由主机(CPU)启动的设备函数但它们是在GPU上执行的,GPU和CPU不通信(除非我们让它们通信)。因此当GPU内核被启动时,CPU将简单地继续运行后续指令,不管它们是启动更多的内核还是执行其他CPU函数。所以如果在内核启动前后分别调用time.time(),则只获得了内核启动所需的时间,而不是计算运行所需的时间。

所以这里就需要进行同步,也就是调用 cuda.synchronize()函数,这个函数将停止主机执行任何其他代码,直到 GPU 完成已在其中启动的每个内核的执行。

为了计算内核执行时间,可以简单地计算内核运行然后同步所需的时间,但是需要使用 time.perf_counter() 或 time.perf_counter_ns() 而不是 time.time()。

在使用 Numba 时,我们还有一个细节需要注意:Numba 是一个 Just-In-Time 编译器,这意味着函数只有在被调用时才会被编译。因此计时函数的第一次调用也会计时编译步骤,这通常要慢得多。所以必须首先通过启动内核然后对其进行同步来编译代码,确保了下一个内核无需编译即可立即运行,这样得到的时间才是准确的。

 from time import perf_counter_ns
 
 # Compile and then clear GPU from tasks
 add_array[blocks_per_grid, threads_per_block](dev_a, dev_b, dev_c)
 cuda.synchronize()
 
 timing = np.empty(101)
 for i in range(timing.size):
     tic = perf_counter_ns()
     add_array[blocks_per_grid, threads_per_block](dev_a, dev_b, dev_c)
     cuda.synchronize()
     toc = perf_counter_ns()
     timing[i] = toc - tic
 timing *= 1e-3  # convert to μs
 
 print(f"Elapsed time: {timing.mean():.0f} ± {timing.std():.0f} μs")
 
 #  Elapsed time: 201 ± 109 μs
 
 # Compile and then clear GPU from tasks
 add_array_gs[blocks_per_grid_gs, threads_per_block](dev_a, dev_b, dev_c)
 cuda.synchronize()
 
 timing_gs = np.empty(101)
 for i in range(timing_gs.size):
     tic = perf_counter_ns()
     add_array_gs[blocks_per_grid_gs, threads_per_block](dev_a, dev_b, dev_c)
     cuda.synchronize()
     toc = perf_counter_ns()
     timing_gs[i] = toc - tic
 timing_gs *= 1e-3  # convert to μs
 
 print(f"Elapsed time: {timing_gs.mean():.0f} ± {timing_gs.std():.0f} μs")
 
 #  Elapsed time: 178 ± 141 μs

对于简单的内核,还可以测量算法的整个过程获得每秒浮点运算的数量。通常以GFLOP/s (giga-FLOP /s)为度量单位。加法操作只包含一个触发器:加法。因此,吞吐量由:

 #              G * FLOP       / timing in s
 gflops    = 1e-9 * dev_a.size * 1e6 / timing.mean()
 gflops_gs = 1e-9 * dev_a.size * 1e6 / timing_gs.mean()
 
 print(f"GFLOP/s (algo 1): {gflops:.2f}")
 print(f"GFLOP/s (algo 2): {gflops_gs:.2f}")
 
 #  GFLOP/s (algo 1): 4.99
 #  GFLOP/s (algo 2): 5.63

一个二维的例子

下面使用一个更复杂的例子,我们创建一个2D内核来对图像应用对数校正。

给定一个值在0到1之间的图像I(x, y),对数校正图像由

Iᵪ(x, y) = γ log₂(1 + I(x, y))

首先让我们获取一些数据!

 import matplotlib.pyplot as plt
 from skimage import data
 
 moon = data.moon().astype(np.float32) / 255.
 
 fig, ax = plt.subplots()
 im = ax.imshow(moon, cmap="gist_earth")
 ax.set_xticks([])
 ax.set_yticks([])
 ax.set_xticklabels([])
 ax.set_yticklabels([])
 fig.colorbar(im)
 fig.tight_layout()

数据在较低端饱和了。几乎没有高于0.6的值。

现在编写核函数。

 import math
 
 # Example 1.5: 2D kernel
 @cuda.jit
 def adjust_log(inp, gain, out):
     ix, iy = cuda.grid(2) # The first index is the fastest dimension
     threads_per_grid_x, threads_per_grid_y = cuda.gridsize(2) #  threads per grid dimension
     
     n0, n1 = inp.shape # The last index is the fastest dimension
     # Stride each dimension independently
     for i0 in range(iy, n0, threads_per_grid_y):
         for i1 in range(ix, n1, threads_per_grid_x):
             out[i0, i1] = gain * math.log2(1 + inp[i0, i1])
 
 threads_per_block_2d = (16, 16)  #  256 threads total
 blocks_per_grid_2d = (64, 64)
 
 moon_gpu = cuda.to_device(moon)
 moon_corr_gpu = cuda.device_array_like(moon_gpu)
 
 adjust_log[blocks_per_grid_2d, threads_per_block_2d](moon_gpu, 1.0, moon_corr_gpu)
 moon_corr = moon_corr_gpu.copy_to_host()
 
 fig, (ax1, ax2) = plt.subplots(1, 2)
 ax1.imshow(moon, cmap="gist_earth")
 ax2.imshow(moon_corr, cmap="gist_earth")
 ax1.set(title="Original image")
 ax2.set(title="Log-corrected image")
 for ax in (ax1, ax2):
     ax.set_xticks([])
     ax.set_yticks([])
     ax.set_xticklabels([])
     ax.set_yticklabels([])
 fig.tight_layout()

第一个for循环从iy开始,第二个最里面的循环从ix开始。为什么选择这个顺序呢?这种选择的内存访问模式更有效。因为第一个网格索引是最快的,所以我们想让它匹配最快的维度:最后一个维度。

结果如下:

总结

本文中介绍了Numba和CUDA的基础知识,我们可以创建简单的CUDA内核,并将其从内存移动到GPU的显存来使用它们。还介绍了如何使用Grid-stride技术在1D和2D数组上迭代。

附录:使用Nvidia的cuda-python来查看设备属性

要对GPU的确切属性进行细粒度控制,可以使用Nvidia提供的较低级别的官方CUDA Python包,代码如下:

 # Need to: pip install --upgrade cuda-python
 
 from cuda.cuda import CUdevice_attribute, cuDeviceGetAttribute, cuDeviceGetName, cuInit
 
 # Initialize CUDA Driver API
 (err,) = cuInit(0)
 
 # Get attributes
 err, DEVICE_NAME = cuDeviceGetName(128, 0)
 DEVICE_NAME = DEVICE_NAME.decode("ascii").replace("\x00", "")
 
 err, MAX_THREADS_PER_BLOCK = cuDeviceGetAttribute(
     CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, 0
 )
 err, MAX_BLOCK_DIM_X = cuDeviceGetAttribute(
     CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, 0
 )
 err, MAX_GRID_DIM_X = cuDeviceGetAttribute(
     CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X, 0
 )
 err, SMs = cuDeviceGetAttribute(
     CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, 0
 )
 
 print(f"Device Name: {DEVICE_NAME}")
 print(f"Maximum number of multiprocessors: {SMs}")
 print(f"Maximum number of threads per block: {MAX_THREADS_PER_BLOCK:10}")
 print(f"Maximum number of blocks per grid:   {MAX_BLOCK_DIM_X:10}")
 print(f"Maximum number of threads per grid:  {MAX_GRID_DIM_X:10}")
 
 #  Device Name: Tesla T4                                                                
 #  Maximum number of multiprocessors: 40
 #  Maximum number of threads per block:       1024
 #  Maximum number of blocks per grid:         1024
 #  Maximum number of threads per grid:  2147483647

作者:Carlos Costa

本文转载于网络 如有侵权请联系删除

相关文章

  • 搭建普罗米修斯Prometheus监控系统「建议收藏」

    大家好,又见面了,我是你们的朋友全栈君。一、普罗米修斯监控概述1、什么是普罗米修斯监控Prometheus(由go语言(golang)开发)是一套开源的监控&报警&时间序列数据库的组合。适合监控docker容器。因为K8S的流行带动了Prometheus的发展。2、官方网站https://prometheus.io/docs/introduction/overview/二、时间序列数据1、什么是时间序列数据时间序列数据(TimeSeriesData):按照时间顺序记录系统、设备状态变化的数据被称为时序数据。应用场景很多,如:无人驾驶车辆中要记录信息、传统证券行业实时交易数据等等。2、基于时间序列数据的特点性能好存储成本低3、普罗米修斯特征多维度数据模型灵活的查询语言不依赖分布式存储,单个服务器节点是自主的以HTTP方式,通过pull模型拉取时间序列数据也可以通过中间网关支持push模型通过服务发现或者静态配置,来发现目标服务对象支持多种多样的图表和界面展示4、普罗米修斯原理架构图三、实验环境准备1、准备服务器grafana服务器(运维成像=>数据转换成图形) pr

  • 6万辆自动驾驶小车将入市!图扑软件构建车联网系统

    《汽车工业蓝皮书:中国商用汽车产业发展报告(2022)》指出,自动驾驶末端配送小车将在未来五年内快速发展,到2025年,中国自动驾驶末端配送市场小车将达到6万辆。7月18日,工业和信息化部装备工业一司组织召开智能网联汽车推进组(ICV-2035)2022年度工作会议。强调,要进一步把握产业变革趋势和发展脉络,持续加强重大问题研究评估,强化各方资源力量协调,加快开展智能网联汽车准入管理试点,推动关键核心技术研发应用,组织更大范围更多场景测试示范,同步推进法规政策完善、技术标准制定、产业生态建设等各项工作,努力推动我国智能网联汽车产业更好更快发展。车联网实现了人们“第二空间”汽车的智能化,同时也是万物智联中的一部分,汽车不再是冰冷的机器,而是有感情、温度的智能硬件。车联网在推动汽车产品升级的同时,数字技术的演进,同样赋予了汽车感知和智慧,让汽车从交通工具向智能终端进化,具有了交互和服务的能力。运用图扑软件自主研发的HTforWeb可视化产品,将车联网上下游数据有效整合于一屏。利用图扑软件提供的多边型呈现地图区块,拼接来实现预期效果,相比之下更加轻量快捷,不需要建模就可以完成,极大降低了实施

  • MySQL权限详解

    设置MySQL用户资源限制通过设置全局变量max_user_connections可以限制所有用户在同一时间连接MySQL实例的数量,但此参数无法对每个用户区别对待,所以MySQL提供了对每个用户的资源限制管理MAX_QUERIES_PER_HOUR:一个用户在一个小时内可以执行查询的次数(基本包含所有语句)MAX_UPDATES_PER_HOUR:一个用户在一个小时内可以执行修改的次数(仅包含修改数据库或表的语句)MAX_CONNECTIONS_PER_HOUR:一个用户在一个小时内可以连接MySQL的时间MAX_USER_CONNECTIONS:一个用户可以在同一时间连接MySQL实例的数量从5.0.3版本开始,对用户‘user’@‘%.example.com’的资源限制是指所有通过example.com域名主机连接user用户的连接,而不是分别指从host1.example.com和host2.example.com主机过来的连接用户资源限制执行操作通过执行createuser/alteruser设置/修改用户的资源限制mysql>CREATEUSER'wsp&#

  • Vivado-hls使用实例

    Vivado-hls使用实例【引言】本系列教程演示如何使用xilinx的HLS工具进行算法的硬件加速。分为三个部分,分别为HLS端IP设计,vivado硬件环境搭建,SDK端软件控制。在HLS端,要将进行硬件加速的软件算法转换为RTL级电路,生成便于嵌入式使用的axi控制端口,进行数据的传输和模块的控制。【HLS介绍】HLS可以将算法直接映射为RTL电路,实现了高层次综合。vivado-HLS可以实现直接使用C,C++以及SystemC语言对Xilinx的FPGA器件进行编程。用户无需手动创建RTL,通过高层次综合生成HDL级的IP核,从而加速IP创建。HLS的官方参考文档主要为:ug871(ug871-vivado-high-level-synthesis-tutorial.pdf)和ug902(ug902-vivado-high-level-synthesis.pdf)。对于VivadoHls来说,输入包括Tesbench,C/C++源代码和Directives,相应的输出为IPCatalog,DSP和SysGen,特别的,一个工程只能有一个顶层函数用于综和,这个顶层函数下面的子函

  • 在mysql中order by是怎样工作的?

    先举一个例子: 创建一张表:CREATETABLE`t`( `id`int(11)NOTNULL, `city`varchar(16)NOTNULL, `name`varchar(16)NOTNULL, `age`int(11)NOTNULL, `addr`varchar(128)DEFAULTNULL, PRIMARYKEY(`id`), KEY`city`(`city`) )ENGINE=InnoDB;复制通过这个下面这段sql进行排序:selectcity,name,agefromtwherecity='杭州'orderbynamelimit1000;复制排序过程:初始化一个sortbuffer我们对city进行了索引的创建所以通过索引将city为杭州的筛选出来;(减少全表扫描)将筛选出来的cityagename字段放在内存中的sortbuffer中(sortbuffer为排序开辟的一块新内存)直到不符合查询的条件。(就算是limit等于1000在这一步也会查出比1000多的数据在这块分页是不起作用的)一直重复第三步将符合条件的在所有数据存入sortbuffe

  • Java Review (十八、面向对象----对象与垃圾回收)

    Java运行时数据区域图一:Java运行时数据区域示意图Java内存运行时区域中的程序计数器、虚拟机栈、本地方法栈随线程而生灭;因此这几个区域的内存分配和回收都具备确定性,不需要过多考虑回收的问题,因为方法结束或者线程结束时,内存自然就跟随着回收了。而Java堆不一样,一个接口中的多个实现类需要的内存可能不一样,一个方法中的多个分支需要的内存也可能不一样,只有在程序处于运行期间时才能知道会创建哪些对象,这部分内存的分配和回收都是动态的,Java的垃圾回收机制所关注的是这部分内存。Java垃圾回收机制具有如下特征:垃圾回收机制只负责回收堆内存中的对象,不会回收任何物理资源(例如数据库连接、网络IO等资源)。程序无法精确控制垃圾回收的运行,垃圾回收会在合适的时候进行。当对象永久性地失去引用后,系统就会在合适的时候回收它所占的内存。在垃圾回收机制回收任何对象之前,总会先调用它的finalize()方法,该方法可能使该对象重新复活(让一个引用变量重新引用该对象),从而导致垃圾回收机制取消回收。对象在内存中的状态当一个对象在堆内存中运行时,根据它被引用变量所引用的状态,可以把它所处的状态分成如下

  • Access重复项查询

    大家好上节介绍了汇总查询,继续介绍选择查询中的重复项查询和不匹配项查询,这两种查询都可以在查询向导中创建,本节主要介绍重复项查询。一、重复项查询重复项查询:将数据库中相同字段的信息内容集合在一起显示,主要用于各种数据的对比分析。在一部分表中,可能会使用自动编号的数据类型ID作为字段的主键,而非使用自然主键。虽然这样也可以保证主键的唯一性,但是记录的数据可能出现重复的情况。此时就可以通过重复项查询,查找出重复项记录,并且可以在查询中将重复的记录删除。二、示例下面复制一个图书表的副本,新建ID字段,数据类型选用自动编号,并作为主键。示例将资本论的信息复制添加最最后一行。下面就利用查询向导中的重复项查询来找出重复项。如下图所示:选择重复查询向导,选择图书表副本。在通过哪些字段选择查找值时,选择书名或者作者名均可。然后选择查询后显示的字段。这里可以全选。指定查询的名称后,选择完成,得到查询的结果。将重复的两条记录显示出来。然后选择其中一个一整行的记录(注意是选择整行),单击鼠标右键,选择删除记录。就可以删除一条记录。(删除记录时要谨慎,删除的是数据库表中的数据。)回到表的数据表视图中,可以看到

  • 基于 Django 信号机制实现类似触发器的效果

    我们都知道,在关系数据库中,为了保证数据完整性,我们都会使用一个叫做触发器的玩意。今天我就基于Django信号机制实现类似触发器的效果,在此之前我先简单介绍一下触发器。触发器触发器(trigger)是数据库提供给程序员和数据分析员来保证数据完整性的一种方法,它是与表事件相关的特殊的存储过程,它的执行不是由程序调用,也不是手工启动,而是由事件来触发,比如当对一个表进行操作(insert,delete,update)时就会激活它执行。触发器经常用于加强数据的完整性约束和业务规则等。一般情况下,常用的触发器总共有6种——增加数据之前的触发器、增加数据之后的触发器、删除数据之前的触发器、删除数据之后的触发器、修改数据之前的触发器、修改数据之后的触发器。触发器的优点触发器可通过数据库中的相关表实现级联更改,不过,通过级联引用完整性约束可以更有效地执行这些更改。触发器可以强制比用CHECK约束定义的约束更为复杂的约束。与CHECK约束不同,触发器可以引用其它表中的列。例如,触发器可以使用另一个表中的SELECT比较插入或更新的数据,以及执行其它操作,如修改数据或显示用户定义错误信息。触发器也可以评

  • 步骤都给你了,还不动手?

    本文标识:L00001 本文编辑:灭霸 编程工具:Linux 阅读时长:5分钟准备镜像和VMwareVMware和rhel-server-6.4-x86_64-dvd.iso这个直接去网上下载安装就好。安装好以后我们会看见一下这个界面建议大家使用英文版的,在企业工作大多数都是英文,先熟悉一下环境。1、首先我们新建一个虚拟机。2、选择典型。3、选择稍后安装操作系统,后期我们自己挂载镜像。4、接下来选择Linux的版本,根据自己下载的镜像版本而定。5、下一步给虚拟机取一个名字,给它一个存放路径,注意:路径最好不要出现中文。6、下面是磁盘大小和,将虚拟磁盘统一成一个文件。7、检查一下如果没有问题点击Finish。8、在左上方会出现一个Linux这就是我们即将安装的虚拟机,然后选择镜像。9、把开始下载的镜像路径添加到上面点击OK10、然后运行此虚拟机。Linux系统安装步骤选择第一项,然后回车:跳过光盘质量测试提示上一步回车后,将出现下面的界面,使用“Tab”键切换到“Skip”,然后回车:点击Next如果鼠标不能移动可用Ctrl+Alt进行切换选择安装过程使用的语言选择安装过程使用的语言:中

  • LeetCode-ZigZag conversion

    ZigzagConversionThestring"PAYPALISHIRING"iswritteninazigzagpatternonagivennumberofrowslikethis:(youmaywanttodisplaythispatterninafixedfontforbetterlegibility) PAHN APLSIIG YIR Andthenreadlinebyline:"PAHNAPLSIIGYIR" Writethecodethatwilltakeastringandmakethisconversiongivenanumberofrows: stringconvert(strings,intnumRows); Example1: Input:s="PAYPALISHIRING",numRows=3 Output:"PAHNAPLSIIGYIR" Example2: Input:s="PAYPALISHIRING",numRows=4 Output:

  • Spring Boot 注册 Servlet 的三种方法,真是太有用了!

    本文栈长教你如何在SpringBoot注册Servlet、Filter、Listener。一、SpringBoot注册SpringBoot提供了ServletRegistrationBean,FilterRegistrationBean,ServletListenerRegistrationBean三个类分别用来注册Servlet,Filter,Listener,下面是Servlet的示例代码。importjavax.servlet.http.HttpServlet; importjavax.servlet.http.HttpServletRequest; importjavax.servlet.http.HttpServletResponse; importjava.io.IOException; /** *@authorJava技术栈 */ publicclassRegisterServletextendsHttpServlet{ @Override protectedvoidservice(HttpServletRequestreq,HttpServletResponsere

  • 博鳌论坛,盖茨和马斯克答李彦宏:我们并不反对人工智能

    特斯拉CEO和CTO埃隆·马斯克(ElonMusk)曾表示,“人工智能如同召唤恶魔”,微软联合创始人也提出要警惕人工智能的观点,这被媒体大肆渲染。然而,在博鳌亚洲论坛上,两位大佬对百度CEO李彦宏明确表示,他们并不反对人工智能。马斯克说,发展人工智能的方向是对的,但是我们不能操之过急。盖茨畅想了未来建立起人工智能的愿景:“希望通过硅来实现,把知识植入进去,包括数学模型,包括因特网,可以把书和知识融入进去,通过这样的基础把人工智能建造起来。”29日上午7:00~8:00,海南博鳌。百度CEO李彦宏主持了一场与微软创始人比尔·盖茨和特斯拉CEO埃隆·马斯克(ElonMusk)的对话,主题是《技术、创新与可持续发展》。以下为对话速记节选:主持人李彦宏:昨天有关人工智能进行了简短的辩论,今天也谈一下。人工智能动的进步带来了很多的争论,百度最近看到在你的访谈当中,你也是说到了一些人工智能的负面,这是有人在访谈当中所说的。比如说火星或者是其他方面的人工智能有负面因素,你对此有什么看法?ElonMusk:实际上我是觉得这是非常不正确的比喻,我并不认为我是火星人。数据智能的风险,并不是说就是超人类,并

  • 学界 | 新型循环神经网络IndRNN:可构建更长更深的RNN(附GitHub实现)

    选自arXiv作者:ShuaiLi等机器之心编译参与:张倩、黄小天近日,澳大利亚伍伦贡大学联合电子科技大学提出一种新型的循环神经网络IndRNN,不仅可以解决传统RNN所存在的梯度消失和梯度爆炸问题,还学习长期依赖关系;此外,借助relu等非饱和激活函数,训练之后IndRNN会变得非常鲁棒,并且通过堆叠多层IndRNN还可以构建比现有RNN更深的网络。实验结果表明,与传统的RNN和LSTM相比,使用IndRNN可以在各种任务中取得更好的结果。同时本文还给出了IndRNN的TensorFlow实现,详见文中GitHub链接。循环神经网络(RNN)[16]已在动作识别[8]、场景标注[4]、语言处理[5]等序列学习问题中获得广泛应用,并且成果显著。与卷积神经网络(CNN)等前馈网络相比,RNN具有循环连接,其中最后的隐藏状态是到下一状态的输入。状态更新可描述如下:其中和分别为时间步t的输入和隐藏状态。、和分别为当前输入的权重、循环输入以及神经元偏差,σ是神经元的逐元素激活函数,N是该RNN层中神经元的数目。由于循环权重矩阵不断相乘,RNN的训练面临着梯度消失和梯度爆炸的问题。长短期记忆(L

  • 汇编14:端口

    端口 在PC系统中,和CPU通过总线相连的芯片除了各种存储器外,还有以下三种芯片: 1、各种接口卡(如显卡、网卡)上的接口芯片,它们控制接口卡进行工作 2、主板上的接口芯片,CPU通过它们对部分外设进行访问 3、其他芯片,用来存储相关的系统信息,或进行相关的输入输出处理 这些芯片中都有一组可以由CPU读写的寄存器,这些寄存器都与CPU的总线相连,且CPU都可以通过控制线发出端口读写命令来对这些寄存器读或写。从CPU的角度来看,可以把这些寄存器都当做端口进行统一编址,建立一个统一的端口地址空间。 CPU可以直接读写的对象有三种:CPU内部的寄存器、内存单元、端口。 端口的读写 访问端口时,端口地址与内存地址一样,通过地址总线来传送,在PC系统中CPU最多可以定位64KB个不同的端口,端口地址的范围为0-65535。 端口读写的指令只有两条:in和out,分别用于从端口读取数据和往端口写入数据。在in和out指令中,只能使用ax或al来存放从端口中读入的数据,或要发送到端口中的数据。访问8位端口时用al,16位用ax。 从60h号端口读入一个字节: inal,60h 复制 执行这条语句实际

  • Revit二次开发——读取shp

    简单记录一下进展: 首先编译shp读写库为X64位,因为本机的Revit2018是64位的。 接着编译C#封装库MapTools,AnyCPU,应该是P-Invoke实现的,有些年头不搞C#了,不知道最新技术是啥。 然后就是在类库项目中调用C#封装库。编写读取shp的代码。 作者:太一吾鱼水 宣言:在此记录自己学习过程中的心得体会,同时积累经验,不断提高自己! 声明:博客写的比较乱,主要是自己看的。如果能对别人有帮助当然更好,不喜勿喷! 文章未经说明均属原创,学习笔记可能有大段的引用,一般会注明参考文献。 欢迎大家留言交流,转载请注明出处。

  • ctrl alt shift +T

    绘制一个图形后,按下ctrl+j,复制图层;接着对新图层按下ctrl+T,出现界定框,按下方向键移动一段距离,回车确认;接着按下ctrl+shift+alt三键不放,并连续按T,每按下一次T,出现一个新的复制图层,并与前俩个的相关比例相同

  • node.js系列(实例):原生node.js实现接收前台post请求提交数据

    技术交流群:821039247 前台界面: 前台代码: 1<formclass="form-horizontal"method="post"action="http:127.0.0.1:3000/post"> 2<divclass="form-group"> 3<labelfor="inputEmail3"class="col-sm-2control-label">姓名</label> 4<divclass="col-sm-4"> 5<inputtype="text"name="username"class="form-control"id="inputEmail3"> 6</div> 7</div> 8<divclass="form-group"> 9<labelclass="col-sm-2control-label">性别</label> 10<divclass="col-sm-4"> 11<labelclass="radio

  • match 和 search 和 indexOf 查找及 正则表达式的 exec 和 test 用法

    functiontest(){ varname="1.087"; varabc="abdwor66kne78xt"; varreg=/\d+/g; log(name.indexOf('.')); log(name.split('.')); log("匹配12a是否有数字:"+reg.test("12a"));//匹配成功返回true失败false log("匹配abc56AUKg789第一次出现的多个数字:"+reg.exec("abc56AUKg789"));//匹配第一个匹配成功的字符串 log("match查找字符中[a-z]的多个字串数组:"+abc.match(/[a-z]+/g));//匹配全部【a-z】字符 log("serach查找字符串中ne出现的位置:"+abc.search(/[ne]+/g));//查找ne的位置 log("indexOf查找字符串中ne出现的位置:"+abc.indexOf('ne'));//查找ne的位置 } functionlog(s){ console.log(s); } 复制      佛语:我

  • 3.12进度报告

    因为之前的爬虫学的不是太好,今天边学爬虫边根据数据中的单位名称在百度地图上爬取对应的地理位置信息,写了一天,本来以为能爬出来了,但是经过一小段循环测试后发现了好多问题,比如有的单位找不到,有的单位爬取位置却不正确。明天继续搞

  • 第十二周总结

    第十二周总结 1.实验源码: packagemain1; importjavax.swing.*; classloginwindow{ privateJFrameframe=newJFrame("登陆窗口");//声明一个窗体对象 privateJButtonsubmit=newJButton("登陆");//声明一个按钮 privateJButtonreset=newJButton("重置");//声明一个按钮 privateJLabelnamelab=newJLabel("用户名:");//声明一个标签 privateJLabelpasslab=newJLabel("密码:");//声明一个标签 privateJLabelinfolab=newJLabel("用户登录系统");//声明一个标签 privateJTextFieldnametest=newJTextField();//定义一个文本域 privateJTextFieldpasstest=newJPasswordField();//定义一个文本域 publicloginwindow(){ frame.setLayou

  • Splay 伸展树 bzoj3224 bzoj3223

    Splay伸展树,平衡树的一种实现方法 splay的精髓在于rotate函数,这里不多作介绍,以及有大牛把原理及实现方法解释的很清楚了,这里只贴一下自己实现的代码,代码参考 :史上最详尽的平衡树(splay)讲解与模板 下面是我自己的代码: bzoj3224  传送门:Tyvj1728普通平衡树 #include"iostream" #include"iomanip" #include"string.h" #include"stack" #include"queue" #include"string" #include"vector" #include"set" #include"map" #include"algorithm" #include"stdio.h" #include"math.h" #pragmacomment(linker,"/STACK:102400000,102400000") #definebug(x)cout<<x<<""<<"UUUUU"<<endl; #definemem(a,x)mem

相关推荐

推荐阅读