1.硬件理解

1.1对应

1.2 不一定是同时执行

• 例如只有13个sm,每个sm有128个core,而我们创建了1百万个threads,就要同步执行
  
• 因此，我们倾向于在block的x维设置为32的倍数,防止浪费warp
  
• warp id打印

#include <stdio.h>
#include <stdlib.h>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

__global__ void print_details_of_warps()
{
	int gid = blockIdx.y * gridDim.x * blockDim.x 
		+ blockIdx.x * blockDim.x + threadIdx.x;

	int warp_id = threadIdx.x / 32;

	int gbid = blockIdx.y * gridDim.x + blockIdx.x;

	printf("tid : %d, bid.x : %d, bid.y : %d, gid : %d, warp_id : %d, gbid : %d \n",
		threadIdx.x, blockIdx.x, blockIdx.y, gid, warp_id, gbid);
}

int main(int argc , char** argv)
{
	dim3 block_size(42);
	dim3 grid_size(2,2);

	print_details_of_warps << <grid_size,block_size >> > ();
	cudaDeviceSynchronize();

	cudaDeviceReset();
	return EXIT_SUCCESS;
}

• 需要注意,条件语句不总是会导致发散,当我申请blocksize为64时会分配2个warp(此处能优化计算速度)
  

• 代码效率计算
  

• grid对应kernal函数
  

• 对kernal函数而言

  • thread-core
  • block-sm
  • grid-device
    
    
    

• 同一个block里执行的数据尽量要靠近
  
  

• CudaDeviceSynchronize (会阻塞CPU，直到所有先前的CUDA调用都完成为止）

• CudaMemcpy（）调用之前会调用CudaDeviceSynchronize

• cudaMemcpyAsync()这个不会调用，用在stream中，pipeline
  
  
  
  
  

• 优化矩阵乘法
  

• 内存中线性分布，一行行的串联