- 博客(15)
- 收藏
- 关注
RSS订阅原创 深度学习算子实现（NEON+CUDA）（持续更新......）
实现一个卷积操作，使用c语言编写的话实际就是四个循环，一个循环遍历输入矩阵的行，一个循环遍历输入矩阵的列，定位到需要卷积的子矩阵，然后一个循环遍历kernel的行，一个循环遍历kernel的列，这两个循环用来计算当前卷积的结果。
2024-01-07 20:30:44
438
原创 GEMM实现（NEON+CUDA）
在Fortran编程当中数组是默认按列存储的，而在C/C++当中通常是按行存储的。A的第一行先和B的第一行按照A的通道为1的元素来进行广播相乘，因为B0就是矩阵B的第一行，按照矩阵的乘法，只能和A的第一行的第一个元素相乘，然后累加到寄存器c0中，这样将B的四个行遍历结束后，就能计算出C矩阵中的一行了。然后将上面的思想用到一张640*640的数组中计算，其实就是对矩阵A的行进行4个为一组的分行块，对矩阵B的列进行分列块，每一次分好的块进行4*4的矩阵计算，就能够迭代出整个矩阵C的结果了。
2023-12-13 22:58:49
1454
原创 CUDA学习笔记（四）
因为在结构体数组中，一个元素内部是一个结构体，存放着不同的数据，那么整个结构体数组在内存中的存放就会是以下排列方式abc abc abc abc，要访问该结构体数组中的全部a元素时，就无法满足合并内存访问，造成带宽浪费。使用固定内存会导致主机端的性能降低，但其实在数据传输的时候使用cudamemcpy时，主机和设备端时同步的，也就是主机端必须等待设别断数据传输完毕才能向下运行，所以尽管主机端性能会降低，但在大数据传输的时候，固定内存还是比较适用的。寄存器：寄存器是距离计算核心最近的内存，读写速度是最快的。
2023-12-07 20:21:51
366
原创 通过例子学习CUDA（一）
其实，在加载数据时，可能并不是加载来的所有字节都是我们需要使用的，有的时候会出现带宽浪费，所以这个指标可以反映我们加载的数据中被用到的比例是否高。该卷积操作使用的是3*3的卷积核，32*32的矩阵。其实从上面四个对比很难看出造成他们之间的性能差距的因素是什么，但很容易看出的是，当block的最内层维度是32的倍数的时候是比较高效的，因为一个warp是32个线程，能够同时加载32个数据，更不容易造成带宽的浪费。该卷积操作时最简单的卷积操作，没有实现padding，步长是1，没有使用共享内存优化。
2023-12-06 19:47:40
332
1
原创 CUDA编程学习笔记（三）
另外，我在网上查询数据后，了解到jetson nano上总共的内存只有4G，我本以为这4G是cpu和gpu共享的内存，从而不需要进行显示的数据传输（使用cudaMalloc，cudaFree，cudaMemcopy等api），但上手操作了一下，发现仍然需要进行显示传输。这就会造成前16个线程先去执行+1操作，而后16个线程等待，等到前16个线程执行完毕后，后16个线程执行-1操作，等到后16个线程执行完毕后，整个warp才调度出去。很明显，jetson nano上只有一个SM，并且只有128个计算核心。
2023-12-01 16:30:15
345
1
原创 CUDA编程学习笔记（二）
gpu上的计算核心是有限的，但我们可以通过线程组织管理来让有限的核心计算超出核心数量的线程，那就是因为需要计算的线程其实是被按照一个一个的warp来进行组织的，而一个warp会被同时执行，当这个warp执行完毕，会通过warp调度来加载下一个warp进行计算，所以gpu能够高效计算超出计算核心数量的线程。在编写程序的时候，会将需要进行并行处理的代码编写为一个kernel，该kernel是送到GPU上给计算核心进行计算的最小单位，一个计算核心同时只能有一个kernel在计算。4.编写简单例子并分析。
2023-11-30 19:28:46
1683
原创 ARM架构NEON intrinsic函数学习笔记（四）
vmlaq_f32：实现128位的，即包含4个float数据的寄存器之间的对应元素的相乘，然后与第三个寄存器相加。vget_high_f32：获取寄存器的高64位包含的两个float数据，vget_low_f32同理；vaddq_f32：实现128位，即包含4个float数据的寄存器之间的对应元素加法操作；vadd_f32：实现64位，即包含2个float数据的寄存器之间的对应元素加法操作；vpadd_f32：实现64位的，即包含2个float数据的寄存器的逐个元素相加；C语言代码（针对循环展开）
2023-11-30 15:58:25
422
原创 ARMv8-A架构寄存器学习篇
sp寄存器：sp是栈寄存器，该寄存器用来存储当前栈空间的栈顶指针，并且sp寄存器可以通过减操作进行栈的开辟，通过加操作进行栈的释放（因为栈空间是从高地址到低地址延伸的，所以减是开辟栈空间，加是释放栈空间）。零寄存器存储的值默认是0，当对零寄存器进行读操作时，读取到的数据是0，当对零寄存器进行写操作时，会忽略写操作，所以零寄存器可以当做是一个只读的寄存器，并且只读的数据是0；使用了wzr零寄存器对栈空间中的地址赋值，因为钙离子中处理的数据是int类型，所以自动使用的是32位的零寄存器。
2023-11-27 14:40:33
448
原创 通过例子学习汇编（二）
这里很肯定的说使用了neon寄存器。一个neon寄存器是128位，处理的数据时int类型，所以一个neon寄存器可以同时处理4个int型数据。所以很明显，开了O3级别优化，编译器自动的使用了neon实现SIMD并行处理数据了。另外，我以为对for循环进行4个一组的循环展开可能会使编译器自动使用neon进行优化，但不开O3的时候编译器并没有自动优化。感觉学习的不多，希望有对这方面了解的大佬和小伙伴可以评论交流啊，互相学习鸭！老实说，开了O3优化后的汇编好多都看不明白呜呜呜。实现一个函数进行前n项和计算。
2023-11-26 14:55:05
362
1
原创 通过例子学习汇编（一）
在开启-O3级别的优化时，最大的变化是func函数没有了调用者局部变量和后续指令的入栈和出栈指令，并且puts函数的调用只是b puts，不是bl puts了，表示func函数调用结束没有返回了。所以说编译器检测到了func和main没有数据依赖，很有可能进行了并行处理了，实现了加速。其中数据的加载分为了两步，先是加载了数据所在的内存页的地址，然后使用偏移量加载该页中的具体数据。arm64中的寄存器由32个，即x0-x31，一个是64bit，也就是8byte，所以一次入栈两个寄存器改变的是16个字节。
2023-11-25 18:09:14
355
原创 ARM架构NEON intrinsic函数学习笔记（三）
rgb图像是三个通道组成的，一个像素点的色彩是由三个通道对应的r，g，b三个颜色来表示的，并且一个像素点的三个r，g，b是连续存储的。对于一次处理16个像素点来讲，图像的像素数量如果不是16的整数倍的话，应该将前面是16整数倍的通过neon进行处理，后面不足16个像素组的像素点只能按照单个uint8的方式处理。另外，由于uint8占8比特，vld3q_u8一次可以将三个128位的寄存器填满，所以在读取rgb数组的时候，一个是可以处理128/8=16个像素点的。
2023-11-25 13:51:44
616
原创 CUDA编程学习笔记（一）
面对cpu具有的强大的逻辑处理能力和gpu强大的数据处理能力，将两者的长处结合起来是很有必要的，所以cuda就是一个将两个架构平台联系在一起的、能够协作编程、协作处理问题的一个平台。而gpu的架构设计完全相反，gpu上搭载的处理核心数大得多，都是上百的数量级，但是没有处理复杂分支预测的逻辑单元。一个核心一次只能运行一个线程，所以多个核心需要计算时，就需要在一个核心上串联地执行，即后续的线程需要等待前面的线程执行完毕才能得到执行。gpu：较多的核心，较强的大数据处理能力，但没有处理复杂逻辑问题的能力。
2023-11-23 19:39:11
368
原创 ARM架构NEON intrinsic函数学习笔记（二）
总结：由于intrinsic函数是neon指令的封装，所以比类似于循环展开这种编译器自动优化的方式效果来得要稍微差一些，但是在加减乘法运算时还是要比不优化要好的。其中可以看出，一次只处理一个数据的效率是最差的，使用neon intrinsic函数次之，最好的是使用了循环展开技巧的情况。首先使用vld1q_f32将元素加载到两个寄存器，然后使用vaddq_f32对两个寄存器当中的数据进行运算（每个寄存器中有128bit数据，共4个float数据），最后使用vst1q_f32将计算结果存储到目标数组。
2023-11-23 15:39:20
98
原创 ARM架构NEON intrinsic函数学习笔记（一）
Neon指令操作的寄存器有两种，一种是128位的寄存器，一种是64位的寄存器（其中64位的寄存器其实就是128位的寄存器没有使用高位的64位而已，其实都是128位的寄存器）。所以以128位寄存器举例的话，一个neon指令，可以simd同时处理16个8bit数据、4个32bit数据、2个64bit数据。其中intrinsic函数是对neon指令的封装，方便使用，连着使用的性能相差不大，但都比手写neon汇编代码带来的性能要差一些。运算操作的最小数据，有8位、16位、32位等。rigister：寄存器。
2023-11-21 20:54:11
130
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人