CUDA编程- - GPU线程的理解 thread,block,grid - 学习记录

GPU线程的理解 thread,block,grid

  • 一、从 cpu 多线程角度理解 gpu 多线程
    • 1、cpu 多线程并行加速
    • 2、gpu多线程并行加速
      • 2.1、cpu 线程与 gpu 线程的理解（核函数）
        • 2.1.1 、第一步：编写核函数
        • 2.1.2、第二步：调用核函数（使用内核函数）
        • 2.1.3、第三步：编写 CMakeLists & 编译代码
  • 二、重要概念 & 与线程索引的直观理解
    • 2.1、重要概念
    • 2.2、dim3与启动内核
    • 2.3、如何找到线程块的索引
    • 2.4、如何找到绝对线程索引
  • 三、参考代码（打印索引）
    • 3.1、打印一维索引
    • 3.2、打印二维索引
    • 3.2、扩展应用 （获取图片坐标）

一、从 cpu 多线程角度理解 gpu 多线程

1、cpu 多线程并行加速

在 cpu 中，用 openmp 并行计算，有限的线程数对 128 进行分组运算。

#pragma omp parallel for
for(int i =0;i<128;i++)
{a[i]=b[i]*c[i];
}

2、gpu多线程并行加速

在 gpu 中，可以直接开启 128 个线程对其进行计算。下面步骤和代码是演示如何开启 128个线程并打印

2.1、cpu 线程与 gpu 线程的理解（核函数）

2.1.1 、第一步：编写核函数
__global__ void some_kernel_func(int *a, int *b, int *c)
{// 初始化线程IDint i = (blockIdx.x * blockDim.x) + threadIdx.x;// 对数组元素进行乘法运算a[i] = b[i] * c[i];// 打印打前处理的进程ID// 可以看到blockIdx并非是按照顺序启动的，这也说明线程块启动的随机性printf("blockIdx.x = %d,blockDimx.x = %d,threadIdx.x = %d\n", blockIdx.x, blockDim.x, threadIdx.x);
}
2.1.2、第二步：调用核函数（使用内核函数）

#）调用语法
kernel_function<<<num_blocks,num_threads>>>(param1,param2,...)

  • num_blocks 线程块，至少保证一个线程块
  • num_threads 执行内核函数的线程数量

#）tips:

1、 some_kernel_func<<<1,128>>>(a,b,c); 调用 some_kernel_func 1*128 次

2、 some_kernel_func<<<2,128>>>(a,b,c); 调用 some_kernel_func 2*128 次

3、如果将 num_blocks 从 1 改成 2 ，则表示 gpu 将启动两倍于之前的线程数量的线程，

在 blockIdx.x = 0 中，i = threadIdx.x
在 blockIdx.x = 1 中， blockDim.x 表示所要求每个线程块启动的线程数量，在这 = 128

2.1.3、第三步：编写 CMakeLists & 编译代码

CMakeLists.txt

cmake_minimum_required(VERSION 2.8 FATAL_ERROR)
project(demo)
add_definitions(-std=c++14)      
find_package(CUDA REQUIRED)# add cuda
include_directories(${CUDA_INCLUDE_DIRS} ) 
message("CUDA_LIBRARIES:${CUDA_LIBRARIES}")
message("CUDA_INCLUDE_DIRS:${CUDA_INCLUDE_DIRS}")
cuda_add_executable(demo print_theardId.cu)# link
target_link_libraries (demo ${CUDA_LIBRARIES})

print_theardId.cu

#include <stdio.h>
#include <stdlib.h>#include <cuda.h>
#include <cuda_runtime.h>/*gpu 中的矩阵乘法*/
__global__ void some_kernel_func(int *a, int *b, int *c)
{// 初始化线程IDint i = (blockIdx.x * blockDim.x) + threadIdx.x;// 对数组元素进行乘法运算a[i] = b[i] * c[i];// 打印打前处理的进程ID// 可以看到blockIdx并非是按照顺序启动的，这也说明线程块启动的随机性printf("blockIdx.x = %d,blockDimx.x = %d,threadIdx.x = %d\n", blockIdx.x, blockDim.x, threadIdx.x);
}int main(void)
{// 初始化指针元素int *a, *b, *c;// 初始化GPU指针元素int *gpu_a, *gpu_b, *gpu_c;// 初始化数组大小int size = 128 * sizeof(int);// 为CPU指针元素分配内存a = (int *)malloc(size);b = (int *)malloc(size);c = (int *)malloc(size);// 为GPU指针元素分配内存cudaMalloc((void **)&gpu_a, size);cudaMalloc((void **)&gpu_b, size);cudaMalloc((void **)&gpu_c, size);// 初始化数组元素for (int i = 0; i < 128; i++){b[i] = i;c[i] = i;}// 将数组元素复制到GPU中cudaMemcpy(gpu_b, b, size, cudaMemcpyHostToDevice);cudaMemcpy(gpu_c, c, size, cudaMemcpyHostToDevice);// 执行GPU核函数some_kernel_func<<<4, 32>>>(gpu_a, gpu_b, gpu_c);// 将GPU中的结果复制到CPU中cudaMemcpy(a, gpu_a, size, cudaMemcpyDeviceToHost);// 释放GPU和CPU中的内存cudaFree(gpu_a);cudaFree(gpu_b);cudaFree(gpu_c);free(a);free(b);free(c);return 0;
}

mkdir build
cd build
cmake ..
make
./demo

部分结果：
在这里插入图片描述

可以看到，
1、gpu 可以直接调用很多个线程，
2、线程数量的多少是由线程块，线程，线程网格等决定的，
3、在核函数中编写单个线程的使用代码，再调用核函数，便可简单的达到 cpu 中 openmp 的多线程方式

二、重要概念 & 与线程索引的直观理解

2.1、重要概念

在这里插入图片描述

gridDim.x – 线程网格X维度的线程块数目
gridDim.y – 线程网格Y维度的线程块数目

blockDim.x – 一个线程块X维度上的线程数量
blockDim.y – 一个线程块Y维度上的线程数量

theadIdx.x – 线程块X维度上的线程数量
theadIdx.y – 线程块Y维度上的线程数量

一般来说：
一个 kernel 对应一个 grid
一个 grid 可以有多个 block，一维~三维
一个 block 可以有多个 thread，一维~三维

2.2、dim3与启动内核

dim3 是CUDA中的特殊数据结构，可用来创建二维的线程块与线程网络
eg：4个线程块，128个线程

dim3 threads_rect(32,4)	// 每个线程块在X方向开启32个线程，Y方向开启4个线程
dim3 blocks_rect(1,4)	//在线程网格上，x方向1个线程块，Y方向4个线程
or
dim3 threads_square(16,8)
dim3 blocks_square(2,2)

以上两种方式线程数都是 32x4=128 ， 16x8=128，只是线程块中线程的排布方式不一样

启动内核
1、 some_kernel_func<<<blocks_rect,threads_rect>>>(a,b,c);

2、 some_kernel_func<<<blocks_square,threads_square>>>(a,b,c);

2.3、如何找到线程块的索引

线程块的索引 x 线程块的大小 + 线程数量的起始点
参考核函数

// 定义ID查询函数
__global__ void what_is_my_id(unsigned int *const block,unsigned int *const thread,unsigned int *const warp,unsigned int *const calc_thread)
{/*线程ID是线程块的索引 x 线程块的大小 + 线程数量的起始点*/const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;block[thread_idx] = blockIdx.x;thread[thread_idx] = threadIdx.x;/*线程束 = 线程ID / 内置变量warpSize*/warp[thread_idx] = thread_idx / warpSize;calc_thread[thread_idx] = thread_idx;
}

来个.cu文件，体验一下这个核函数，// 编译方法同上

#include <stdio.h>
#include <stdlib.h>#include "cuda.h"
#include "cuda_runtime.h"// 定义ID查询函数
__global__ void what_is_my_id(unsigned int *const block,unsigned int *const thread,unsigned int *const warp,unsigned int *const calc_thread)
{/*线程ID是线程块的索引 x 线程块的大小 + 线程数量的起始点*/const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;block[thread_idx] = blockIdx.x;thread[thread_idx] = threadIdx.x;/*线程束 = 线程ID / 内置变量warpSize*/warp[thread_idx] = thread_idx / warpSize;calc_thread[thread_idx] = thread_idx;
}// 定义数组大小
#define ARRAY_SIZE 1024
// 定义数组字节大小
#define ARRAY_BYTES ARRAY_SIZE * sizeof(unsigned int)// 声明主机下参数
unsigned int cpu_block[ARRAY_SIZE];
unsigned int cpu_thread[ARRAY_SIZE];
unsigned int cpu_warp[ARRAY_SIZE];
unsigned int cpu_calc_thread[ARRAY_SIZE];// 定义主函数
int main(void)
{// 总线程数量为 2 x 64 = 128// 初始化线程块和线程数量const unsigned int num_blocks = 2;const unsigned int num_threads = 64;char ch;// 声明设备下参数unsigned int *gpu_block, *gpu_thread, *gpu_warp, *gpu_calc_thread;// 声明循环数量unsigned int i;// 为设备下参数分配内存cudaMalloc((void **)&gpu_block, ARRAY_BYTES);cudaMalloc((void **)&gpu_thread, ARRAY_BYTES);cudaMalloc((void **)&gpu_warp, ARRAY_BYTES);cudaMalloc((void **)&gpu_calc_thread, ARRAY_BYTES);// 调用核函数what_is_my_id<<<num_blocks, num_threads>>>(gpu_block, gpu_thread, gpu_warp, gpu_calc_thread);// 将设备下参数复制到主机下cudaMemcpy(cpu_block, gpu_block, ARRAY_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_thread, gpu_thread, ARRAY_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_warp, gpu_warp, ARRAY_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_calc_thread, gpu_calc_thread, ARRAY_BYTES, cudaMemcpyDeviceToHost);// 释放GPU内存cudaFree(gpu_block);cudaFree(gpu_thread);cudaFree(gpu_warp);cudaFree(gpu_calc_thread);// 循环打印结果for (i = 0; i < ARRAY_SIZE; i++){printf("Calculated Thread: %d - Block: %d - Warp: %d - Thread: %d\n", cpu_calc_thread[i], cpu_block[i], cpu_warp[i], cpu_thread[i]);}return 0;
}

2.4、如何找到绝对线程索引

thread_idx = ( (gridDim.x * blockDim.x ) * idy ) + idx;

绝对线程索引 = 当前行索引 * 每行线程总数 + x方向的偏移
参考核函数

/*定义线程id计算函数*/
__global__ void what_is_my_id_2d_A(unsigned int *const block_x,unsigned int *const block_y,unsigned int *const thread,unsigned int *const calc_thread,unsigned int *const x_thread,unsigned int *const y_thread,unsigned int *const grid_dimx,unsigned int *const block_dimx,unsigned int *const grid_dimy,unsigned int *const block_dimy)
{/*获得线程索引*/const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;/*计算线程id计算公式：线程ID = ((网格维度x * 块维度x) * 线程idy) + 线程idx(作为x维度上的偏移)*/const unsigned int thread_idx = ((gridDim.x * blockDim.x) * idy) + idx;/*获取线程块的索引*/block_x[thread_idx] = blockIdx.x;block_y[thread_idx] = blockIdx.y;/*获取线程的索引*/thread[thread_idx] = threadIdx.x;/*计算线程id*/calc_thread[thread_idx] = thread_idx;/*获取线程的x维度索引*/x_thread[thread_idx] = idx;/*获取线程的y维度索引*/y_thread[thread_idx] = idy;/*获取网格维度的X，Y值*/grid_dimx[thread_idx] = gridDim.x;grid_dimy[thread_idx] = gridDim.y;/*获取block_dimy*/block_dimx[thread_idx] = blockDim.x;
}

来个.cu文件，体验一下这个核函数，// 编译方法同上

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>/*定义线程id计算函数*/
__global__ void what_is_my_id_2d_A(unsigned int *const block_x,unsigned int *const block_y,unsigned int *const thread,unsigned int *const calc_thread,unsigned int *const x_thread,unsigned int *const y_thread,unsigned int *const grid_dimx,unsigned int *const block_dimx,unsigned int *const grid_dimy,unsigned int *const block_dimy)
{/*获得线程索引*/const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;/*计算线程id计算公式：线程ID = ((网格维度x * 块维度x) * 线程idy) + 线程idx(作为x维度上的偏移)*/const unsigned int thread_idx = ((gridDim.x * blockDim.x) * idy) + idx;/*获取线程块的索引*/block_x[thread_idx] = blockIdx.x;block_y[thread_idx] = blockIdx.y;/*获取线程的索引*/thread[thread_idx] = threadIdx.x;/*计算线程id*/calc_thread[thread_idx] = thread_idx;/*获取线程的x维度索引*/x_thread[thread_idx] = idx;/*获取线程的y维度索引*/y_thread[thread_idx] = idy;/*获取网格维度的X，Y值*/grid_dimx[thread_idx] = gridDim.x;grid_dimy[thread_idx] = gridDim.y;/*获取block_dimy*/block_dimx[thread_idx] = blockDim.x;
}/*定义矩阵宽度以及大小*/
#define ARRAY_SIZE_X 32
#define ARRAY_SIZE_Y 16
#define ARRAY_SIZE_IN_BYTES (ARRAY_SIZE_X * ARRAY_SIZE_Y * sizeof(unsigned int))/*声明CPU端上的各项参数内存*/
unsigned int *cpu_block_x[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_block_y[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_warp[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_calc_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_x_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_y_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_grid_dimx[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_grid_dimy[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_block_dimx[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int *cpu_block_dimy[ARRAY_SIZE_Y][ARRAY_SIZE_X];int main(void)
{const dim3 thread_rect = (32, 4);/*注意这里的块的dim3值为1x4*/const dim3 block_rect = (1, 4);/*初始化矩形线程分布启动项*/const dim3 thread_square = (16, 8);/*注意这里的块的dim3值为2x2*/const dim3 block_square = (2, 2);/*定义一个临时指针用于打印信息*/char ch;/*定义GPU端上的各项参数内存*/unsigned int *gpu_block_x;unsigned int *gpu_block_y;unsigned int *gpu_thread;unsigned int *gpu_warp;unsigned int *gpu_calc_thread;unsigned int *gpu_x_thread;unsigned int *gpu_y_thread;unsigned int *gpu_grid_dimx;unsigned int *gpu_grid_dimy;unsigned int *gpu_block_dimx;/*分配GPU端上的各项参数内存*/cudaMalloc((void **)&gpu_block_x, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_block_y, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_thread, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_warp, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_calc_thread, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_x_thread, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_y_thread, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_grid_dimx, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_grid_dimy, ARRAY_SIZE_IN_BYTES);cudaMalloc((void **)&gpu_block_dimx, ARRAY_SIZE_IN_BYTES);/*调用核函数*/for (int kernel = 0; kernel < 2; kernel++){switch (kernel){case 0:/*执行矩形配置核函数*/what_is_my_id_2d_A<<<block_rect, thread_rect>>>(gpu_block_x, gpu_block_y, gpu_thread, gpu_warp, gpu_calc_thread, gpu_x_thread, gpu_y_thread, gpu_grid_dimx, gpu_grid_dimy, gpu_block_dimx);break;case 1:/*执行方形配置核函数*/what_is_my_id_2d_A<<<block_square, thread_square>>>(gpu_block_x, gpu_block_y, gpu_thread, gpu_warp, gpu_calc_thread, gpu_x_thread, gpu_y_thread, gpu_grid_dimx, gpu_grid_dimy, gpu_block_dimx);break;default:exit(1);break;}/*将GPU端上的各项参数内存拷贝到CPU端上*/cudaMemcpy(cpu_block_x, gpu_block_x, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_block_y, gpu_block_y, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_thread, gpu_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_warp, gpu_warp, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_calc_thread, gpu_calc_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_x_thread, gpu_x_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_y_thread, gpu_y_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_grid_dimx, gpu_grid_dimx, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_grid_dimy, gpu_grid_dimy, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);cudaMemcpy(cpu_block_dimx, gpu_block_dimx, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);printf("\n kernel %d\n", kernel);/*打印结果*/for (int y = 0; y < ARRAY_SIZE_Y; y++){for (int x = 0; x < ARRAY_SIZE_X; x++){printf("CT: %2u Bkx: %1u TID: %2u YTID: %2u XTID: %2u GDX: %1u BDX: %1u GDY: %1u BDY:%1U\n", cpu_calc_thread[y * ARRAY_SIZE_X + x], cpu_block_x[y * ARRAY_SIZE_X + x], cpu_thread[y * ARRAY_SIZE_X + x], cpu_y_thread[y * ARRAY_SIZE_X + x], cpu_x_thread[y * ARRAY_SIZE_X + x], cpu_grid_dimx[y * ARRAY_SIZE_X + x], cpu_block_dimx[y * ARRAY_SIZE_X + x], cpu_grid_dimy[y * ARRAY_SIZE_X + x], cpu_block_y[y * ARRAY_SIZE_X + x]);}/*每行打印完后按任意键继续*/ch = getchar();}printf("Press any key to continue\n");ch = getchar();}/*释放GPU端上的各项参数内存*/cudaFree(gpu_block_x);cudaFree(gpu_block_y);cudaFree(gpu_thread);cudaFree(gpu_warp);cudaFree(gpu_calc_thread);cudaFree(gpu_x_thread);cudaFree(gpu_y_thread);cudaFree(gpu_grid_dimx);cudaFree(gpu_grid_dimy);cudaFree(gpu_block_dimx);
}

其中有个代码片段

   const dim3 thread_rect = (32, 4);/*注意这里的块的dim3值为1x4*/const dim3 block_rect = (1, 4);const dim3 thread_square = (16, 8);/*注意这里的块的dim3值为2x2*/const dim3 block_square = (2, 2);

如图理解，都是 2x2 / 1x4 = 四个线程块；每一块 32x4 / 16x8 =128个线程。这是两种不同的线程块布局方式。

但是一般会选择长方形的布局方式。
1、要以行的方式进行连续访问内存，而不是列的方式
2、同一个线程块可以通过共享内存进行通信
3、同一个线程束中的线程存储访问合并在一起了，长方形布局只需要一次访问操作就可以获得连续的内存数据 // 正方形要两次访问

在这里插入图片描述

三、参考代码（打印索引）

3.1、打印一维索引

场景：
一个数组有 8 个数据，要开 8 个线程去访问。
我们想切成 2 个block 访问，所以一个 blcok 就有 4 个线程

在这里插入图片描述
所以 线程设置如下：一个 block里面4个线程，一个grid里面2个block

一维索引的设置如下：

    dim3 block(4);// 一个 block 里面 4 个线程dim3 grid(2);//	一个 grid 里面 2 个 block	

#include <cuda_runtime.h>
#include <stdio.h>__global__ void print_idx_kernel(){printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,threadIdx.z, threadIdx.y, threadIdx.x);
}__global__ void print_dim_kernel(){printf("grid dimension: (%3d, %3d, %3d), block dimension: (%3d, %3d, %3d)\n",gridDim.z, gridDim.y, gridDim.x,blockDim.z, blockDim.y, blockDim.x);
}__global__ void print_thread_idx_per_block_kernel(){int index = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;printf("block idx: (%3d, %3d, %3d), thread idx: %3d\n",blockIdx.z, blockIdx.y, blockIdx.x,index);
}__global__ void print_thread_idx_per_grid_kernel(){int bSize  = blockDim.z * blockDim.y * blockDim.x;int bIndex = blockIdx.z * gridDim.x * gridDim.y + \blockIdx.y * gridDim.x + \blockIdx.x;int tIndex = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;int index  = bIndex * bSize + tIndex;printf("block idx: %3d, thread idx in block: %3d, thread idx: %3d\n", bIndex, tIndex, index);
}void print_one_dim(){int inputSize = 8;int blockDim = 4;int gridDim = inputSize / blockDim;dim3 block(blockDim);//4dim3 grid(gridDim);//2print_idx_kernel<<<grid, block>>>();//print_dim_kernel<<<grid, block>>>();//print_thread_idx_per_block_kernel<<<grid, block>>>();//print_thread_idx_per_grid_kernel<<<grid, block>>>();cudaDeviceSynchronize();	//用于同步
}int main() {print_one_dim();return 0;
}

核函数及其结果：

  • 8个线程，8个输出；
  • 索引都是从 z到y到x的；

1、线程块与线程

__global__ void print_idx_kernel(){printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,threadIdx.z, threadIdx.y, threadIdx.x);
}

在这里插入图片描述

2、线程网格与线程块

__global__ void print_dim_kernel(){printf("grid dimension: (%3d, %3d, %3d), block dimension: (%3d, %3d, %3d)\n",gridDim.z, gridDim.y, gridDim.x,blockDim.z, blockDim.y, blockDim.x);
}

1x1x2=2
1x1x4=4

在这里插入图片描述

3、在 block 里面寻找每个线程的索引

__global__ void print_thread_idx_per_block_kernel(){int index = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;printf("block idx: (%3d, %3d, %3d), thread idx: %3d\n",blockIdx.z, blockIdx.y, blockIdx.x,index);
}

可以根据下面的图来理解访问顺序：

在这里插入图片描述

结果：
在这里插入图片描述
4、在 grid 里面寻找每个线程索引

__global__ void print_thread_idx_per_grid_kernel(){int bSize  = blockDim.z * blockDim.y * blockDim.x;	// block 的线程大小int bIndex = blockIdx.z * gridDim.x * gridDim.y + \blockIdx.y * gridDim.x + \blockIdx.x;int tIndex = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;int index  = bIndex * bSize + tIndex;printf("block idx: %3d, thread idx in block: %3d, thread idx: %3d\n", bIndex, tIndex, index);
}

可以根据下面的图来理解访问顺序：实际上就是从一堆方块里面找到那个红点

在这里插入图片描述
结果：（thread 从 0 ~ 7 ）
在这里插入图片描述

3.2、打印二维索引

#include <cuda_runtime.h>
#include <stdio.h>__global__ void print_idx_kernel(){printf("block idx: (%3d, %3d, %3d), thread idx: (%3d, %3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,threadIdx.z, threadIdx.y, threadIdx.x);
}__global__ void print_dim_kernel(){printf("grid dimension: (%3d, %3d, %3d), block dimension: (%3d, %3d, %3d)\n",gridDim.z, gridDim.y, gridDim.x,blockDim.z, blockDim.y, blockDim.x);
}__global__ void print_thread_idx_per_block_kernel(){int index = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;printf("block idx: (%3d, %3d, %3d), thread idx: %3d\n",blockIdx.z, blockIdx.y, blockIdx.x,index);
}__global__ void print_thread_idx_per_grid_kernel(){int bSize  = blockDim.z * blockDim.y * blockDim.x;int bIndex = blockIdx.z * gridDim.x * gridDim.y + \blockIdx.y * gridDim.x + \blockIdx.x;int tIndex = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;int index  = bIndex * bSize + tIndex;printf("block idx: %3d, thread idx in block: %3d, thread idx: %3d\n", bIndex, tIndex, index);
}void print_two_dim(){int inputWidth = 4;int blockDim = 2;int gridDim = inputWidth / blockDim;dim3 block(blockDim, blockDim);// 2 , 2dim3 grid(gridDim, gridDim);    //2,2print_idx_kernel<<<grid, block>>>();// print_dim_kernel<<<grid, block>>>();// print_thread_idx_per_block_kernel<<<grid, block>>>();//print_thread_idx_per_grid_kernel<<<grid, block>>>();cudaDeviceSynchronize();
}int main() {print_two_dim();return 0;
}

3.2、扩展应用 （获取图片坐标）

原理其实就是同上面（在 grid 里面寻找每个线程索引）一样，这里为了方便看，再次贴一次图。

在这里插入图片描述
重写一个核函数，比之前的方便看

__global__ void print_cord_kernel(){int index = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;int x  = blockIdx.x * blockDim.x + threadIdx.x;int y  = blockIdx.y * blockDim.y + threadIdx.y;printf("block idx: (%3d, %3d, %3d), thread idx: %3d, cord: (%3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,index, x, y);
}

完整的 .cu 文件如下：

#include <cuda_runtime.h>
#include <stdio.h>__global__ void print_cord_kernel(){int index = threadIdx.z * blockDim.x * blockDim.y + \threadIdx.y * blockDim.x + \threadIdx.x;int x  = blockIdx.x * blockDim.x + threadIdx.x;int y  = blockIdx.y * blockDim.y + threadIdx.y;printf("block idx: (%3d, %3d, %3d), thread idx: %3d, cord: (%3d, %3d)\n",blockIdx.z, blockIdx.y, blockIdx.x,index, x, y);
}void print_cord(){int inputWidth = 4;int blockDim = 2;int gridDim = inputWidth / blockDim;dim3 block(blockDim, blockDim);dim3 grid(gridDim, gridDim);print_cord_kernel<<<grid, block>>>();cudaDeviceSynchronize();
}int main() {print_cord();return 0;
}

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/658314.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

网络代理用途

网络代理用途

网络代理的用途广泛&#xff0c;常用于代理爬虫&#xff0c;代理VPN &#xff0c;代理注入等。使用网络 代理能够将入侵痕迹进一步减少&#xff0c;能够突破自身IP的访问限制&#xff0c;提高访问速度&#xff0c; 以及隐藏真实IP &#xff0c;还能起到一定的防止攻击的作用。下…
阅读更多...
ROS方向第二次汇报（2）

ROS方向第二次汇报（2）

文章目录 1.本方向内学习内容&#xff1a;1.1.动作&#xff1a;1.1.1.案例接口定义:1.1.2.案例通信模型&#xff1a;1.1.3.服务器端代码&#xff1a;1.1.4.客户端源代码&#xff1a;1.1.5.动作命令行操作&#xff1a; 1.2.参数&#xff1a;1.2.1.查看参数列表&#xff1a;1.2.2…
阅读更多...
dvwa靶场xss储存型

dvwa靶场xss储存型

xss储存型 xxs储存型lowmessage框插入恶意代码name栏插入恶意代码 medium绕过方法 high xxs储存型 攻击者事先将恶意代码上传或储存到漏洞服务器中&#xff0c;只要受害者浏览包含此恶意代码的页面就会执行恶意代码。产生层面:后端漏洞特征:持久性的、前端执行、储存在后端数据…
阅读更多...
4JS表达式和运算符expression and operator

4JS表达式和运算符expression and operator

表达式&#xff08;expression&#xff09;JavaScript中的一个短语&#xff0c;JavaScript解释器会将其计算&#xff08;evaluate&#xff09;出一个结果。程序中的常量是最简单的一类表达式。变量名也是一种简单的表达式&#xff0c;它的值就是赋值给变量的值。复杂表达式是由…
阅读更多...
类与对象（中篇）

类与对象（中篇）

1、类的6个默认成员函数 如果一个类中什么成员都没有，简称为空类。 空类中真的什么都没有吗？并不是，任何类在什么都不写时，编译器会自动生成以下6个默认成员函数。 默认成员函数：用户没有显式实现，编译器会生成的成员函数称为默认成员函数。 2、构造函数---初始…
阅读更多...
Java技术栈 —— Spring MVC 与 Spring Boot

Java技术栈 —— Spring MVC 与 Spring Boot

参考文章或视频链接[1] Spring vs. Spring Boot vs. Spring MVC[2] Key Differences Between Spring vs Spring Boot vs Spring MVC
阅读更多...
CRM系统主要干什么的

CRM系统主要干什么的

阅读本文&#xff0c;你将了解&#xff1a;一、CRM系统是什么&#xff1b;二、CRM系统主要干什么的&#xff1b;三、CRM系统在企业管理中的重要作用&#xff1b;四、企业落地案例分享——大吉包装。 本文所提及的功能演示和图片内容均来自于我们公司正在使用的简道云CRM系统&a…
阅读更多...
【JVM】类加载流程

【JVM】类加载流程

目录 1.加载 2.链接 &#xff08;1&#xff09;校验 &#xff08;2&#xff09;准备 &#xff08;3&#xff09;解析 3.初始化 4.使用 5.卸载 1.加载 加载阶段&#xff0c;简言之&#xff0c;查找并加载类的二进制数据&#xff0c;生成 Class 的实例 在加载类时&#x…
阅读更多...
【C++航海王：追寻罗杰的编程之路】引用、内联、auto关键字、基于范围的for、指针空值nullptr

【C++航海王：追寻罗杰的编程之路】引用、内联、auto关键字、基于范围的for、指针空值nullptr

目录 1 -> 引用 1.1 -> 引用概念 1.2 -> 引用特性 1.3 -> 常引用 1.4 -> 使用场景 1.5 -> 传值、传引用效率比较 1.6 -> 值和引用作为返回值类型的性能比较 1.7 -> 引用和指针的区别 2 -> 内联函数 2.1 -> 概念 2.2 -> 特性 3 -…
阅读更多...
Linux ---- Shell编程三剑客之AWK

Linux ---- Shell编程三剑客之AWK

一、awk处理文本工具 1、awk概述 awk 是一种处理文本文件的语言&#xff0c;是一个强大的文本分析工具。AWK是专门为文本处理设计的编程语言&#xff0c;也是行处理软件&#xff0c;通常用于扫描、过滤、统计汇总工作。用来处理列。数据可以来自标准输入也可以是管道或文件。…
阅读更多...
共享的IP隔一段时间就变？用这种方法可以不需要知道电脑IP

共享的IP隔一段时间就变？用这种方法可以不需要知道电脑IP

前言 一般来说，电脑接入路由器之后，IP是由路由器自动分配的（DHCP），但如果隔一段时间不开机连接路由器，或者更换了别的网卡进行连接，自动分配的IP就会更改。 比如你手机连接着电脑的共享IP：192.168.1.10，但过段时间之后，电脑的IP突然变成了192.168.1.11，那么你的所有…
阅读更多...
Qt+css绘制标题

Qt+css绘制标题

之前学过html和小程序&#xff0c;帮老师做项目的时候也用过vue&#xff0c;在想qt绘制界面是不是也可以使用css,然后查了一些资料&#xff0c;绘制了一个标题&#xff0c;准备用到智能家居的上位机上面。 成果 源码 重写了paintEvent函数和TimeEvent函数&#xff0c;一个用于绘…
阅读更多...
幻兽帕鲁服务器多少钱一个？26元，阿里云腾讯云华为云

幻兽帕鲁服务器多少钱一个？26元，阿里云腾讯云华为云

2024年幻兽帕鲁服务器价格表更新&#xff0c;阿里云、腾讯云和华为云Palworld服务器报价大全&#xff0c;4核16G幻兽帕鲁专用服务器阿里云26元、腾讯云32元、华为云26元&#xff0c;阿腾云atengyun.com分享幻兽帕鲁服务器优惠价格表&#xff0c;多配置报价&#xff1a; 幻兽帕鲁…
阅读更多...
透明拼接屏造型：多样拼接与影响因素

透明拼接屏造型：多样拼接与影响因素

透明拼接屏&#xff0c;以其独特的透明显示效果和灵活的拼接方式&#xff0c;在现代显示领域中独树一帜。其造型多样&#xff0c;包括横屏拼接、竖屏拼接、异形拼接以及定制拼接等多种方式&#xff0c;满足了不同场景和应用的需求。尼伽小编将详细介绍这些拼接方式&#xff0c;…
阅读更多...
PR转场模板|超级炫酷故障特效电影游戏视频转场PR模板剪辑素材

PR转场模板|超级炫酷故障特效电影游戏视频转场PR模板剪辑素材

premiere转场&#xff0c;包含200个带有Sound FX的独特视频转场效果。加强剪辑视频视觉效果&#xff0c;在镜头之间的剪辑和添加文字动画&#xff01; MYFX Extension可帮助您一键浏览和应用预设&#xff01;可以喜爱预设&#xff0c;并拥有自己亲手挑选的库。如果您有任何问题…
阅读更多...
字符数组的学习

字符数组的学习

前言&#xff1a; 在前面我们介绍过字符型数据是以字符的ASCII码储存在存储单元中&#xff0c;一般占一个字节&#xff0c;由于 ASCII码也属于整数类型&#xff0c;因此在C99标准中把字符类型归纳为整数类型中的一种&#xff0c;由于字符数据 的应用比较广泛&#xff0c;尤其…
阅读更多...
抽象类（Java）、模板方法设计模式

抽象类（Java）、模板方法设计模式

一、概念 在Java中有abstract关键字&#xff0c;就是抽象的意思&#xff0c;可用来修饰类和成员方法。 用abstract来修饰类&#xff0c;那这个类就是抽象类&#xff1b;修饰方法&#xff0c;那这个方法就是抽象方法。 修饰符 abstract class 类名{修饰符 abstract 返回值类型…
阅读更多...
【数据结构 02】队列

【数据结构 02】队列

一、原理 队列通常是链表结构&#xff0c;只允许在一端进行数据插入&#xff0c;在另一端进行数据删除。 队列的特性是链式存储&#xff08;随机增删&#xff09;和先进先出&#xff08;FIFO&#xff1a;First In First Out&#xff09;。 队列的缺陷&#xff1a; 不支持随机…
阅读更多...
看员工聊天记录，监控员工电脑聊天记录软件有哪些?

看员工聊天记录，监控员工电脑聊天记录软件有哪些?

企业监控员工电脑聊天记录软件是指企业为了管理员工、提高工作效率和保护公司机密而采取的一种技术手段。随着互联网的发展和普及&#xff0c;员工在工作时间内使用聊天工具进行沟通已经成为常态&#xff0c;因此企业需要一种有效的方式来监控和管理员工的聊天记录。 为什么要监…
阅读更多...
【算法与数据结构】198、213、337LeetCode打家劫舍I, II, III

【算法与数据结构】198、213、337LeetCode打家劫舍I, II, III

文章目录 一、198、打家劫舍二、213、打家劫舍 II三、337、打家劫舍III三、完整代码 所有的LeetCode题解索引&#xff0c;可以看这篇文章——【算法和数据结构】LeetCode题解。 一、198、打家劫舍 思路分析&#xff1a;打家劫舍是动态规划的的经典题目。本题的难点在于递归公式…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023