在GPU上开发大规模并行应用程序时,需要一个调试器,GDB调试器能够处理系统中每个GPU上同时运行的数千个线程。CUDA-GDB提供了无缝的调试体验,可以同时调试应用程序的CPU和GPU部分。 就像GDB一样,CUDA-GDB提供了基于控制台的调试界面,可以从本地系统或具有Telnet或SSH访问权限的任何远程系统的命令行中使用。如果您更喜欢使用GUI前端进行调试,则CUDA-GDB还支持与DDD,EMACS或Nsight Eclipse Edition集成 。 CUDA-GDB是用于调试在Linux和QNX上运行的CUDA应用程序的NVIDIA工具。CUDA-GDB是GNU项目调试器GDB的扩展。该工具为开发人员提供了一种调试实际硬件上运行的CUDA应用程序的机制。这使开发人员可以调试应用程序,而不会出现模拟和仿真环境带来的潜在变化。 CUDA-GDB在Linux上运行,并针对Linux和QNX系统。 CUDA-GDB旨在为使用者提供一个无缝调试环境,该环境允许在同一应用程序中同时调试GPU和CPU代码。正如CUDA C中的编程是C编程的扩展一样,使用CUDA-GDB进行调试是使用GDB进行调试的自然扩展。现有的GDB调试功能固有地用于调试主机代码,并且还提供了其他功能来支持调试CUDA设备代码。 CUDA-GDB支持调试C / C ++和Fortran CUDA应用程序。(Fortran调试支持仅限于64位Linux操作系统)CUDA-GDB可以调试NVCC编译器支持的所有C ++功能。 CUDA-GDB允许用户为单步CUDA应用程序设置断点,还可以检查和修改硬件上运行的任何给定线程的内存和变量。 CUDA-GDB支持调试所有CUDA应用程序,无论它们使用CUDA驱动程序API,CUDA运行时API还是同时使用这两者。 CUDA-GDB支持调试内核,这些内核已针对特定的CUDA体系结构进行了编译,例如 sm_75 要么 sm_80,但还支持调试在运行时编译的内核,称为即时编译或简称JIT编译。 这里插一下什么是QNX的系统 http://www.gnu.org/software/ddd/ http://www.gnu.org/software/emacs/ https://developer.nvidia.com/nsight-eclipse-edition 以上这些都是调试器的前端,就是一个好看的壳子。 https://docs.nvidia.com/cuda/cuda-gdb/index.html 以及有使用的文档 https://docs.nvidia.com/cuda/cuda-quick-start-guide/#ubuntu-x86_64 https://docs.nvidia.com/cuda/cuda-installation-guide-linux/ 我们的Nano是ARM64的Ubuntu的系统 这里先看一下GPU的硬件 https://developer.nvidia.com/zh-cn/cuda-gpus 我们的Nano的算力是这样的 yunswj@yunswj-desktop:~$ uname -m && cat /etc/*release aarch64 DISTRIB_ID=Ubuntu DISTRIB_RELEASE=18.04 DISTRIB_CODENAME=bionic DISTRIB_DESCRIPTION="Ubuntu 18.04.5 LTS" # R32 (release), REVISION: 4.4, GCID: 23942405, BOARD: t210ref, EABI: aarch64, DATE: Fri Oct 16 19:44:43 UTC 2020 NAME="Ubuntu" VERSION="18.04.5 LTS (Bionic Beaver)" ID=ubuntu ID_LIKE=debian PRETTY_NAME="Ubuntu 18.04.5 LTS" VERSION_ID="18.04" HOME_URL="https://www./" SUPPORT_URL="https://help./" BUG_REPORT_URL="https://bugs./ubuntu/" PRIVACY_POLICY_URL="https://www./legal/terms-and-policies/privacy-policy" VERSION_CODENAME=bionic UBUNTU_CODENAME=bionic sudo nano ~/.bashrc export CUBA_HOME=/usr/local/cuda-10.0 export LD_LIBRARY_PATH=/usr/local/cuda-10.0/lib64:$LD_LIBRARY_PATH export PATH=/usr/local/cuda-10.0/bin:$PATH Ctrl+X
然后打y source ~/.bashrc nvcc -v nvcc fatal:No input files specified; use option --help for more information https://docs.nvidia.com/cuda/cuda-installation-guide-linux/ 在这里是环境变量的设置的文档,目前我的系统里面已经有了,我就不重头安装了 在tmp里面是编译的中间的缓存文件区 gdb是调试工具 我可以找到它的路径在哪里 https://github.com/inducer/pycuda/releases 此时我想安装一下cudapy 将下载的cudapy文件拉进机器 在这里 验证 解压 tar zxvf pycuda-2019.1.2.tar.gz cd pycuda-2019.1.2/ python3 configure.py --cuda-root=/usr/local/cuda-10.2 sudo python3 setup.py install import numpy as np import pycuda.autoinit import pycuda.driver as cuda from pycuda.compiler import SourceModule mod = SourceModule(""" #define BLOCK_SIZE 16
typedef struct { int width; int height; int stride; int __padding; //为了和64位的elements指针对齐 float* elements; } Matrix;
// 读取矩阵元素 __device__ float GetElement(const Matrix A, int row, int col) { return A.elements[row * A.stride + col]; }
// 赋值矩阵元素 __device__ void SetElement(Matrix A, int row, int col, float value) { A.elements[row * A.stride + col] = value; }
// 获取 16x16 的子矩阵 __device__ Matrix GetSubMatrix(Matrix A, int row, int col) { Matrix Asub; Asub.width = BLOCK_SIZE; Asub.height = BLOCK_SIZE; Asub.stride = A.stride; Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col]; return Asub; }
__global__ void matrix_mul(Matrix *A, Matrix *B, Matrix *C) { int blockRow = blockIdx.y; int blockCol = blockIdx.x; int row = threadIdx.y; int col = threadIdx.x;
Matrix Csub = GetSubMatrix(*C, blockRow, blockCol);
// 每个线程通过累加Cvalue计算Csub的一个值 float Cvalue = 0;
// 为了计算Csub遍历所有需要的Asub和Bsub for (int m = 0; m < (A->width / BLOCK_SIZE); ++m) { Matrix Asub = GetSubMatrix(*A, blockRow, m); Matrix Bsub = GetSubMatrix(*B, m, blockCol); __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; As[row][col] = GetElement(Asub, row, col); Bs[row][col] = GetElement(Bsub, row, col);
__syncthreads(); for (int e = 0; e < BLOCK_SIZE; ++e) Cvalue += As[row][e] * Bs[e][col];
__syncthreads(); }
SetElement(Csub, row, col, Cvalue); } """)
class MatrixStruct(object): def __init__(self, array): self._cptr = None
self.shape, self.dtype = array.shape, array.dtype self.width = np.int32(self.shape[1]) self.height = np.int32(self.shape[0]) self.stride = self.width self.elements = cuda.to_device(array) # 分配内存并拷贝数组数据至device,返回其地址
def send_to_gpu(self): self._cptr = cuda.mem_alloc(self.nbytes()) # 分配一个C结构体所占的内存 cuda.memcpy_htod(int(self._cptr), self.width.tobytes()) # 拷贝数据至device,下同 cuda.memcpy_htod(int(self._cptr)+4, self.height.tobytes()) cuda.memcpy_htod(int(self._cptr)+8, self.stride.tobytes()) cuda.memcpy_htod(int(self._cptr)+16, np.intp(int(self.elements)).tobytes())
def get_from_gpu(self): return cuda.from_device(self.elements, self.shape, self.dtype) # 从device取回数组数据 def nbytes(self): return self.width.nbytes * 4 + np.intp(0).nbytes
a = np.random.randn(400,400).astype(np.float32) b = np.random.randn(400,400).astype(np.float32) c = np.zeros_like(a)
A = MatrixStruct(a) B = MatrixStruct(b) C = MatrixStruct(c) A.send_to_gpu() B.send_to_gpu() C.send_to_gpu()
matrix_mul = mod.get_function("matrix_mul") matrix_mul(A._cptr, B._cptr, C._cptr, block=(16,16,1), grid=(25,25)) result = C.get_from_gpu() print(np.dot(a,b)) print(result) 最后写一个矩阵程序验证 之后将调试
|
|