#include <stdio.h> #include <cuda_runtime.h> /* CUDA核函数 __global__ 核函数的前缀定义 - 使用__global__修饰的函数,必须是void无返回值 - __global__核函数修饰,必须是nvcc编译才有效,否则无效 - __global__修饰的函数,使用name<<<grid, block, memory, stream>>>(params)启动核函数 - 启动在host,但是执行在device __device__ 修饰的函数,只能够在设备上执行,设备上调用(例如核函数内调用) - nvidia提供了很多很多内置设备函数,比如日常的cos、sin之类的 - 在nvidia团队中,不同内置函数的api接口版本号,被称之为,计算能力 定义如下: __device__,函数执行在设备上 __global__,函数执行在设备上,但是调用在host上。定义核函数的符号 __host__,函数执行在host上,调用也在host上 */ // sigmoid 不能够使用 sigmoid<<<1, 3>>>这种启动他 // 你也不能够直接 sigmoid(0.1) // 只能够在核函数内调用他 __device__ float sigmoid(float value){ return 1 / (1 + exp(-value)); } __global__ void compute(float* a, float* b, float* c){ /* 线程layout的概念,启动的线程会被设计为grid和block,如同提供的参数那样 这个layout的概念是虚拟的,通过cuda驱动实现真实硬件映射,抽象了一个中间层(调度层) - 如果我们有4352 Core - 如果我们需要启动 5000 个线程 - 抽象层他会把5000个线程安排到各个Core中去执行,根据情况来执行数次 - 每次调度单位WarpSize,如果启动的线程不足,也会执行WarpSize,不过core是非激活状态而已 你想告诉他启动多少个线程,通过gridDim和blockDim告诉他 线程数 = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z 2147483647 是 有符号整数(int)的最大值 65535 是 无符号短整数(int16, short)的最大值 gridDim的最大值范围: (x,y,z): (2147483647, 65535, 65535) blockDim的最大值范围:(x,y,z): (1024, 1024, 64) 他的定义,其实在 device_launch_parameters.h uint3 __device_builtin__ __STORAGE__ threadIdx; uint3 __device_builtin__ __STORAGE__ blockIdx; dim3 __device_builtin__ __STORAGE__ blockDim; dim3 __device_builtin__ __STORAGE__ gridDim; int __device_builtin__ __STORAGE__ warpSize; 获取线程ID,进行数据操作 数据索引,是通过blockIdx和threadIdx计算得到 gridDim告诉你Grid的大小 ,blockDim告诉你block大小 blockIdx告诉你所在Grid内的索引, threadIdx告诉你所在block的索引 把gridDim和blockDim设想为一个tensor gridDim的shape = gridDim.z x gridDim.y x gridDim.x blockDim的shape = blockDim.z x blockDim.y x blockDim.x 最终的启动线程的shape维度是:gridDim.z x gridDim.y x gridDim.x x blockDim.z x blockDim.y x blockDim.x 如果启动的线程是个6维度tensor,那么索引,是不是也可以类似 blockIdx.z, blockIdx.y, blockIdx.x, threadIdx.z, threadIdx.y, threadIdx.x 仅仅是在这个场景需要把6个维度索引变为连续的内存单元索引 如果我们有6个维度a, b, c, d, e, f,如果我们有6个位置索引u, v, w, x, y, z position = ((((u * b + v) * c + w) * d + x) * e + y) * f + z */ // gridDim = 1 x 1 x 1 // blockDim = 3 x 1 x 1 // int position = blockDim.x * blockIdx.x + threadIdx.x; c[position] = a[position] * sigmoid(b[position]); } int main(){ // 定义3个数组,host指针 const int num = 3; float a[num] = {1, 2, 3}; float b[num] = {5, 7, 9}; float c[num] = {0}; // 定义3个设备指针,device指针 size_t size_array = sizeof(c); float* device_a = nullptr; float* device_b = nullptr; float* device_c = nullptr; // 分配设备空间,大小是size_array,单位byte cudaMalloc(&device_a, size_array); cudaMalloc(&device_b, size_array); cudaMalloc(&device_c, size_array); // 把数据从host复制到device,其实就是主机复制到显卡 // 复制的是a和b cudaMemcpy(device_a, a, size_array, cudaMemcpyHostToDevice); cudaMemcpy(device_b, b, size_array, cudaMemcpyHostToDevice); // 执行核函数,结果放到c上 compute<<<dim3(1), dim3(3)>>>(device_a, device_b, device_c); // 把计算后的结果c复制回主机 cudaMemcpy(c, device_c, size_array, cudaMemcpyDeviceToHost); // 查看主机上的c内容是多少 for(int i = 0; i < num; ++i) printf("c[%d] = %f\n", i, c[i]); return 0; }