cuda_1

此帖为记录cuda学习笔记,可能有错误,恳请各位指出,谢谢。
如果用GPU解决我们的问题首先要把CPU上的数据复制到GPU上,通过PCIe或NVLink总线。
第二步就是实际执行 ,例如说运行cuda内核。这就是完成我们的计算任务了。
第三部就是计算完成后从GPU复制到CPU,通过网络或磁盘传输。
image

一些语法:

__global__ void mykernel(void){
    }

为了让这段代码转换为可用的GPU代码,需要加一个装饰器,也就是__global__,__global__向编译器发出信号,表明这是一个需要编译的函数,以便它能在GPU上运行。(就是告诉GPU,这个是需要运行在GPU上的函数,但是是来自其他涉笔额 )
GPU的编译器驱动就是NVCC,它调用多个编译器和其他工具想脚本一样编译代码,同时它还可以把代码划分为宿主和设备两部分。host上的代码可能会通过gcc/g++编译。通过__global__定义后 全局函数中的代码编译成可在gpu上运行的形式。

调用核函数 :mykernel函数在GPU上开始执行,还传递了内核启动配置参数1,1

mykernel<<<1,1>>>( );

同时GPU也类似CPU有内存分配复制的函数:

cudaMalloc(), cudaFree(), cudaMemcpy()

这些API利用指针来引用内存空间或者定义内存分配。
需要注意的是CPU memory的指针不要在 device code中解引用,GPU的指针也不要在host code中解引用。
这是两个不同处理器的两个不同memory的分配。
在cuda中,add<<<1,1>>>( )和add<<<N,1>>>( )的意思是不一样的,
这是将内核启动n个块,第二个参数1实际上指的是线程。我们告诉cuda,我们启动了n个块,每个块包含1个线程。所有这些n个块都能够在某个程序上并行执行。线程与块的统称称为gird。
在CUDA中grid的子集是block,一个block代表一个或一组worke(比如一个add( )。
然后一组block称为grid,每个block可以用索引表示blockIdx.x,例如:

__global__ void add(int *a, int *b, int *c) {
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

add参数列表包含三个指针 a,b,c .这三个指针用于引用三个向量A B C。
函数体内,返回类型为void。
blockIdx.x:这是一个内置变量,表示当前线程块在网格中的索引。每个线程块有一个唯一的 blockIdx.x 值。
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];:这行代码表示,对应于 blockIdx.x 的每个线程块执行一个向量加法操作,即将数组 a 和数组 b 中对应位置的元素相加,并将结果存储在数组 c 的相同位置。
而对于host code来说:

#define N 512
int main(void){
	int *a,*b,*c;  //host cpoies of a,b,c
	int *d_a,*d_b,*d_c;  //device copies of a,b,c
	int size = N * sizeof(int);
	//cuda分配空间给d_a ,大小为size void**为指向指针的指针
	cudaMalloc((void **) &d_a, size);
	cudaMalloc((void**)&d_b,size);
	cudaMalloc(void **)&d_c,size;
	//在主机内存中分配一个包含N个整数的数组,并将指针存到a中
	//用整数随机初始化数组a
	a = (int *)malloc(size);  random_ints(a,N);
	b = (int *)malloc(size); random_ints(b,N);
	c = (int *)malloc(size); 
	//把数据copy到device中
	cudaMemcpy(d_a,a,size,cudaMemcpyHostToDevie);
	cudaMemcpy(d_b,b,size,cudaMemcpyHostToDevie);
	//启动add()kernel在GPU上N个blocks
add<<<N,1>>>(d_a,d_b,d_c);

//把计算的数据写回到host
cudaMemcpy(c,d_c,size,cudaMemcpyDeviceToHost);
//释放
free(a); free(b); free(c);
cudaFree(d _a);cudaFree(d_b);cuda_Free(d_c);
;}

了解了这些,我们对cuda有了个宏观的概念,最大的是grid,其次是block,在最后是thread。thread存在于工作层次的最底层。所以我们可以把代码修改为:

__global__ void add(int *a, int *b, int *c) {
c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}

但是想要激发thread的并行性,我们必须要修改内核启动参数的第二个数字add<<<1,N>>>( );

所以综合上述两个例子,我们第一次使用了单线程的N个block和1个block的N个线程来进行并行向量加法。
如果同时使用N个线程N个block呢?

在这之前我们先看下数据索引

image
可以看到每个block都是唯一的,但是thread却不是唯一的,在每个block中都有thread。所以这些无法满足全局索引。那么为了满足全局索引,

其实这就跟多维数组一样,需要查找的时候要用全局索引,所以同多维数组的索引值:

int index = threadIdx.x + blockIdx.x *  blockDim.x;

image

如果面对多维呢

#define N (2048*2048)
#define THREADS_PER_BLOCK 512
int main(void) {
int *a, *b, *c; // host copies of a, b, c
int *d_a, *d_b, *d_c; // device copies of a, b, c
int size = N * sizeof(int);
// Alloc space for device copies of a, b, c
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_b, size);
cudaMalloc((void **)&d_c, size);
// Alloc space for host copies of a, b, c and setup input values
a = (int *)malloc(size); random_ints(a, N);
b = (int *)malloc(size); random_ints(b, N);
c = (int *)malloc(size);
/ Copy inputs to device
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
// Launch add() kernel on GPU
add<<<N/THREADS_PER_BLOCK,THREADS_PER_BLOCK>>>(d_a, d_b, d_c);
// Copy result back to host
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
// Cleanup
free(a); free(b); free(c);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0;
}

我们可以看到唯一和一维数组的代码不同的是add<<< >>>中的数字。并行计算了N/512个块,每个块中有512个thread。
同时我们还需处理异常情况,例如溢出:

__global__ void add(int *a, int *b, int *c, int n) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
if (index < n)
c[index] = a[index] + b[index];
}

这里的索引在计算前先来个if判断防止溢出,确保需要的线程进行运算。
同时还需要更新kernel:

add<<<(N + M-1) / M,M>>>(d_a, d_b, d_c, N);

假设我们有 N 个元素,每个块包含 M 个线程:

(N + M - 1):这是为了确保当 N 不能被 M 整除时,可以正确地计算需要多少个块。通过加上 M - 1,可以向上取整。例如,如果 N = 10 和 M = 4,那么我们需要 3 个块(10 元素分成 3 个块,每块 4 个线程,但最后一个块只有 2 个线程)。
/ M:这是计算总共需要多少个块。将总线程数除以每个块的线程数,就得到了块的数量。

通常情况下会启动超过所需数量的thread,而那些而外的thread由内核代码if进行处理

在cuda中多维线程:
多维线程
多维网格和多维线程本质还是一维的,GPU物理上不分块。上图就是一个二维的线程,注意第一个block是0,0 第二个block是1,0 第三个block是2,0 第四个block是0,1 这是因为block的变化首先是以x变化的,thread也是同理,都是以x先变化的。不同于普通的矩阵

所以对于二维线程来说:

int tid = threadIdx.y * blockDim.x + threadIdx.x;
int bid = blockIdx.y * gridDim.x + blockIdx.x;

上面两行代码是在二维线程中的全局索引,可以看到结尾都是加上threadIdx.x,这也印证了我们刚才说的 是以x先变化的。


下面是二维网格二维线程块的具体描述

二维网格二维线程块
三维网格三维线程块的具体描述
在这里插入图片描述注意blockDim.x都是已经给定的值就是一个block中多少列,blockDim.y就是有多少行,blockDim.z就是有多少层。类似数据结构中计算数组的那个全局索引是类似的。
所以常见的一维索引如下:

//一维grid 一维block
//blockIdx.x就是在第几个block中 threadIdx.x就是在第几个线程
int blockId = blockIdx.x;
int id = blockIdx.x*blockDim.x + threadIdx.x
//一维grid 二维block
int blockId = blockIdx.x;
//前部分是之前block的总线程数,第二部分是当前block对应的线程数。
int id = blockIdx.x * (blockDim.x * blockDim.y) + threadIdx.x + threadIdx.y*blockDim.x;

//一维grid 三维block
int blockIdx = blockIdx.x;
int id = blockIdx.x * blockDim.x * blockDim.y * blockDim.z + blockDim.x *blockDim.y* threadIdx.z + blockIdx.y * blockDim.x     + threadIdx.x;
//二维grid 一维block
int blockid = blockIdx.y * gridDim.x + blockidx.x;
int id = blockid * blockDim.x  + threadIdx.x;

//二维grid 二维block
int blockid = blockIdx.y * gridDim.x + blockIdx.x;
int id = blockIdx * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;

//三维grid 三维block
int blockid = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
int id = blockid * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + theadIdx.x;





1、定义核函数:用 global 关键字定义一个核函数。核函数的返回类型必须是 void。

__global__ void kernel_function(int *a, int *b, int *c) {
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  c[index] = a[index] + b[index];
}

2、调用核函数:在主机代码(host code)中,使用特殊的语法调用核函数,指定执行配置(execution configuration),包括网格(grid)和块(block)的尺寸。

int main() {
  int *d_a, *d_b, *d_c;
  int size = N * sizeof(int);
  
  // 分配设备内存
  cudaMalloc((void**)&d_a, size);
  cudaMalloc((void**)&d_b, size);
  cudaMalloc((void**)&d_c, size);
  // 拷贝数据到设备
  cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
  // 执行核函数
  int blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
  kernel_function<<<blocks, threadsPerBlock>>>(d_a, d_b, d_c);
  // 拷贝结果回主机
  cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
  // 释放设备内存
  cudaFree(d_a);
  cudaFree(d_b);
  cudaFree(d_c);    
  return 0;
}

所以总结成一个模板来说,就可以简单的认为如下流程:

#include

__global_void 函数名(arg...){
	***kernel content
}
int main(void){
设置GPU设备
分配host device内存
初始化主机数据
数据从主机复制到设备
调用核函数在设备中进行计算
将计算得到的数据从设备传给主机
释放主机与设备内存
}	


解释:if(CUDA_FOUND) message(STATUS "Found CUDA Toolkit v${CUDA_VERSION_STRING}") enable_language(CUDA) set(HAVE_CUDA TRUE) if (CMAKE_CUDA_COMPILER_ID STREQUAL "NVIDIA") if(${CUDA_VERSION_STRING} VERSION_GREATER_EQUAL "11.1") execute_process(COMMAND ${CMAKE_CUDA_COMPILER} --list-gpu-code RESULT_VARIABLE EXIT_CODE OUTPUT_VARIABLE OUTPUT_VAL) if(EXIT_CODE EQUAL 0) #Remove sm_ string(REPLACE "sm_" "" OUTPUT_VAL ${OUTPUT_VAL}) #Convert to list string(REPLACE "\n" ";" __CUDA_ARCH_BIN ${OUTPUT_VAL}) #Remove last empty entry list(REMOVE_AT __CUDA_ARCH_BIN -1) else() message(FATAL_ERROR "Failed to run NVCC to get list of GPU codes: ${EXIT_CODE}") endif() elseif(${CUDA_VERSION_STRING} VERSION_GREATER_EQUAL "11.0") set(__CUDA_ARCH_BIN "35;37;50;52;53;60;61;62;70;72;75;80") elseif(${CUDA_VERSION_STRING} VERSION_GREATER_EQUAL "10.0") set(__CUDA_ARCH_BIN "30;32;35;37;50;52;53;60;61;62;70;72;75") elseif(${CUDA_VERSION_STRING} VERSION_GREATER_EQUAL "9.1") set(__CUDA_ARCH_BIN "30;32;35;37;50;52;53;60;61;62;70;72") else() set(__CUDA_ARCH_BIN "30;32;35;37;50;52;53;60;61;62;70") endif() else() message(FATAL_ERROR "Unsupported CUDA compiler ${CMAKE_CUDA_COMPILER_ID}.") endif() set(CUDA_ARCH_BIN ${__CUDA_ARCH_BIN} CACHE STRING "Specify 'real' GPU architectures to build binaries for") if(POLICY CMP0104) cmake_policy(SET CMP0104 NEW) set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCH_BIN}) message(STATUS "CMAKE_CUDA_ARCHITECTURES: ${CMAKE_CUDA_ARCHITECTURES}") #Add empty project as its not required with newer CMake add_library(pcl_cuda INTERFACE) else() # Generate SASS set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCH_BIN}) # Generate PTX for last architecture list(GET CUDA_ARCH_BIN -1 ver) set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode arch=compute_${ver},code=compute_${ver}") message(STATUS "CMAKE_CUDA_FLAGS: ${CMAKE_CUDA_FLAGS}") add_library(pcl_cuda INTERFACE) target_include_directories(pcl_cuda INTERFACE ${CUDA_TOOLKIT_INCLUDE}) endif () endif()
05-30
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值