最近发现了一条真理，那就是科研项目中遇到问题，千万别企图绕过去，因为，最终还是会发现，那个悬而未决的问题会最终把你带回原地。废话不多说，先盗用大佬的一张图，在CUDA架构下, 显示芯片执行时的最小单位是thread. 数个thread可以组成一个block. 一个block中的thread能存取同一块共享的内存(shared memory), 而且可以快速进行同步的动作, 特别要注意, 这是块(block)同步. 不同block中的thread无法存取同一个共享的内存, 因此无法直接互通或进行同步. 因此, 不同block中的thread能合作的程度是比较低的)然后依据thread, block和grid, 有着不同的存储. 核心就是thread. 可以结合下图进行理解:

它的理念是把矩阵分成小块，一个线程计算出的Csub就是对应一个元素，然后你想象所有线程是并行的，所以所有的元素计算就可以理解过来了，但是单纯的一个线程是解决不了这个问题的，因为它小循环迭代的时候是需要所有元素的内容的，这也是使用__syncthreads()的作用，而你在使用global memory的时候，不需要用这个，因为每个线程各做各的，都需要去读global memory,这也是使用shared memory的作用，它可以减少对global 的访问。
CPU上矩阵乘法代码，在GPU中，核函数的的ID操作方式也是基于这样的循环去操作线程ID的，从而有助于矩阵乘法在GPU中并行执行。根据之前的矩阵乘法实验结果来看，矩阵数目越大，发而能把GPU的优良性能发挥出来，随着矩阵数量的增多，时间的减少也是一个量级一个量级的减少。两个1000×1000的矩阵相称，时间是0.221341，两个100000×100000的矩阵相称，时间是0.000025，这个结果很惊讶了。1000000×1000000时间是0.160396，
/* CPUMatMultiply：CPU下矩阵乘法
 * a:第一个矩阵指针，表示a[M][N]
 * b:第二个矩阵指针，表示b[N][S]
 * result:结果矩阵，表示为result[M][S]
 */
 
#include <iostream>
using namespace std;
 
void CPUMatMultiply(int M, int N, int S, float* a, float* b, float* c)
{
    for (int Row = 0; Row < M; ++ Row)
		for (int Col = 0; Col < N; ++ Col )
		{
			float sum = 0;
			for (int i = 0; i < N; ++i)
			{
				float A = a[Row * N + i];
				float B = b[Col + i * S];
				sum += A * B;
			}
			c[Row* S + Col] = sum;
		}
}
 
int main()
{
	//这里将矩阵按照行优先转换成了一维的形式
	float a[6] = { 11.4, 24, 33.5,   45, 55 ,32.4 }; //2×3的矩阵
	float b[12] = {12,43,22.4, 31.3,  12,324,23,12,  44.4,23.4,65.3,73};//3×4的矩阵
	float c[8] = { 0 }; //2×4的结果矩阵
	int M=2,N=3,S=4;
	//上面两个矩阵比较简单，但是我实际上需要操作的矩阵是很多的， a矩阵大约是3×3的矩阵，而b矩阵大到3×100000，所以目前我的问题是如何将这样大的矩阵赋值给指针a,b,c.
 
	MatrixMulOnHost(M, N, S, a, b, c);
 
	//输出结果
    for (int i=0;i<8;i++)
	{
		cout<<c[i]<<endl;
	}
 
    return 0;
}

通过操作线程来执行矩阵的乘法，如果矩阵乘法的核函数写的如同下述 for循环里的，意味着该线程下的数据是与矩阵的数据对应的。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <iostream>
using namespace std;
 
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
 
__global__ void MatrixMulKernle(int m, int n, int k, float *A,float  *B, float *C)
{
	int Row=blockIdx.y*blockDim.y+threadIdx.y;
	int Col=blockIdx.x*blockDim.x+threadIdx.x;
 
	if((Row<m) && (Col<k))
	{
		float Cvalue=0.0;
		for(int i=0;i<n;++i)
		     Cvalue+=A[Row*n+i]*B[Col+i*k];
		C[Row*k+Col]=Cvalue;
	}
}

下述代码块是矩阵乘法核函数的另外一种写法，其中，blockIdx、blockDim、threadIdx、gridDim都是CUDA的内置变量。blockIdx、threadIdx：分别CUDA中线程块的ID、线程的ID。blockDim、gridDim：分别是CUDA中线程块的维度，线程网格的维度。if语句是为了判断是否是对矩阵中的数据进行操作，防止在当线程threadId超过了M*S时，使用result[threadId]出现越界行为。后续的操作则是矩阵乘法的简单实现。 但是我感觉下面这种核函数中的for循环写的不如上面的for循环写的清晰，上面的for  循环将线程对应到了相应的矩阵元素，而下面的for循环只是模糊的匹配了线程下的值。

https://blog.csdn.net/dcrmg/article/details/54867507该篇博客是关于不同grid，block划分下的线程计算公式。按照和函数的公式对应，grid与block都是二维的。
/* gpuMatMultKernel：GPU下矩阵乘法核函数
*  a:第一个矩阵指针，表示a[M][N]
*  b:第二个矩阵指针，表示b[N][S]
*  result:结果矩阵，表示result[M][S]
*/
__global__ void gpuMatMultKernel(const int *a, const int *b, int *result, const int M, const int N, const int S)
{
    int threadId = (blockIdx.y * blockDim.y + threadIdx.y) * gridDim.x * blockDim.x 
                    + blockIdx.x * blockDim.x + threadIdx.x;
    if (threadId < M * S)
    {
        int row = threadId % S;
        int column = threadId % m;

        result[threadId] = 0;
        for (int i = 0; i < N; i++)
        {
            result[threadId] += a[row * N + i] * b[i * S + column];
        }
    }
}

下述的代码块是实现两个矩阵相加的核函数，与矩阵相乘不同的是，这块的线程划分是一维的。通过操作线程，使得相应线程下的数据相加。

```cpp
#include <stdio.h>

__global__ void add(float * x, float *y, float * z, int n){
        int index = threadIdx.x + blockIdx.x * blockDim.x;
        int stride = blockDim.x * gridDim.x;
        
        for (int i = index; i < n; i += stride){
                z[i] = x[i] + y[i];
        }
}

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <iostream>
using namespace std;
#define TILE_WIDTH 16  
 
__global__ void MatrixMulKernle(int m, int n, int k, float *A,float  *B, float *C)
{
	 //申请共享内存，存在于每个block中
	__shared__ float ds_A[TILE_WIDTH][TILE_WIDTH]; 
	__shared__ float ds_B[TILE_WIDTH][TILE_WIDTH];
 
	//简化坐标记法,出现下面6个表示的地方就是并行的地方。
	int bx = blockIdx.x;		int by = blockIdx.y;
	int tx = threadIdx.x;		int ty = threadIdx.y;
 
	//确定结果矩阵中的行和列
	int Row = by * TILE_WIDTH + ty;
	int Col = bx * TILE_WIDTH + tx;
 
	//临时变量
	float Cvalue = 0;
 
	//循环读入A,B瓦片，计算结果矩阵，分阶段进行计算
	for (int t=0; t<(n-1)/TILE_WIDTH+1; ++t)
	{
		//将A,B矩阵瓦片化的结果放入shared memory中，每个线程加载相应于C元素的A/B矩阵元素
		if (Row < m && t * TILE_WIDTH + tx < n)		//越界处理，满足任意大小的矩阵相乘（可选）
			//ds_A[tx][ty] = A[t*TILE_WIDTH + tx][Row];
		    ds_A[tx][ty] = A[Row*n+t*TILE_WIDTH+tx];//以合并的方式加载瓦片
		else
			ds_A[tx][ty] = 0.0;
 
		if (t * TILE_WIDTH + ty < n && Col < k)
			//ds_B[tx][ty] = B[Col][t*TILE_WIDTH + ty];
            ds_B[tx][ty] = B[(t*TILE_WIDTH + ty)*k+Col];
		else
			ds_B[tx][ty] = 0.0;	
 
		//保证tile中所有的元素被加载
		__syncthreads();
		
		for (int i = 0; i < TILE_WIDTH; ++i)
            Cvalue += ds_A[i][ty] * ds_B[tx][i];//从shared memory中取值
 
		//确保所有线程完成计算后，进行下一个阶段的计算
		__syncthreads();
 
		if(Row < m && Col < k)
			C[Row*k+Col]=Cvalue;		
	}
}

/* gpuMatMultWithSharedKernel：GPU下使用shared内存的矩阵乘法
*  a:第一个矩阵指针，表示a[height_A][width_A]
*  b:第二个矩阵指针，表示b[width_A][width_B]
*  result:结果矩阵，表示result[height_A][width_B]
*/
template<int BLOCK_SIZE>
__global__ void gpuMatMultWithSharedKernel(const int *a, const int *b, int *result, const int height_A, const int width_A, const int width_B)
{
//block index
    int block_x = blockIdx.x;
    int block_y = blockIdx.y;
// Thread index
    int thread_x = threadIdx.x;
    int thread_y = threadIdx.y;

    if ((thread_y + block_y * blockDim.y) * width_B + block_x * blockDim.x + thread_x >= height_A * width_B)
    {
        return;
    }
    // Index of the first sub-matrix of A processed by the block
    //这是A的第一个细分矩阵的索引,利用block处理（by=blockIdx.y）
    //也是从这里让aBegin根据不同block进行处理
    //不同block拥有不同的aBegin，即不同的block线程下的循环初始值的不同
    const int begin_a = block_y * BLOCK_SIZE * width_A;
    // Index of the last sub-matrix of A processed by the block
    //这是A的最后一个细分矩阵的索引，利用block处理
    //注意wA是A的width 
    const int end_a = begin_a + width_A - 1;
    // Step size used to iterate through the sub-matrices of A
    //A的细分矩阵的步距
    const int step_a = BLOCK_SIZE;
    // Index of the first sub-matrix of B processed by the block
    //这是B的第一个细分矩阵的索引，也是利用block处理（bx=blockIdx.x）
    const int begin_b = block_x * BLOCK_SIZE;
    // Step size used to iterate through the sub-matrices of B
    //B的细分矩阵的迭代用步距
    const int step_b = BLOCK_SIZE * width_B;

    int result_temp = 0;

    for (int index_a = begin_a, int index_b = begin_b;
        index_a < end_a; index_a += step_a, index_b += step_b)
    {  
        // Declaration of the shared memory array SubMat_A used to
        // store the sub-matrix of A为一个block即A的sub-matrix分配一个As[32][32]的共享内存
        __shared__ int SubMat_A[BLOCK_SIZE][BLOCK_SIZE];
        // Declaration of the shared memory array SubMat_B used to
        // store the sub-matrix of B为一个block即B的sub-matrix分配一个Bs[32][32]的共享内存
        __shared__ int SubMat_B[BLOCK_SIZE][BLOCK_SIZE];

        // Load the matrices from device memory to shared memory; each thread loads
        // one element of each matrix从GPU的内存加载矩阵到恭喜那个内存。每个线程加载一个元素
        SubMat_A[thread_y][thread_x] = a[index_a + thread_y * width_A + thread_x];
        SubMat_B[thread_y][thread_x] = b[index_b + thread_y * width_B + thread_x];
        
        // Synchronize to make sure the matrices are loaded
        __syncthreads();

        for (int i = 0; i < BLOCK_SIZE; i++)
        {
            result_temp += SubMat_A[thread_y][i] * SubMat_B[i][thread_x];
        }

        __syncthreads();
    }
    // Write the block sub-matrix to device memory;
    // each thread writes one element

    int begin_result = block_y * BLOCK_SIZE * width_B + begin_b;
    result[begin_result + thread_y * width_B + thread_x] = result_temp;
}

上述代码来自https://github.com/NVIDIA/cuda-samples，很好的学习例子。上述核函数来自与matrixMul矩阵乘法的并行运算，每次计算矩阵的一块数据。利用共享内存的共享功能，每次将一块数据保存到共享内存中使得一个线程块同时调用数据进行计算当前块相对应得矩阵乘法结果值。

代码 shared int SubMat_A中的__shared__声明变量SubMat_A为共享内存中保存的变量。然后将数组中的数据提取到变量SubMat_A中保存在共享内存。

__syncthreads()对线程块中的线程进行同步，确保对__shared__进行下面的操作时上面的操作已经完成。

两个for循环完成了当前线程块对应矩阵子块的乘法结果计算。
下一次我写关于托管内存以及数据在 host 与device之间传递的区别。

	

一、使用批量消息提升服务端处理能力
    批处理是提升系统吞吐量的有效方法。在Kafka内部，消息都是以“批”为单位处理的。
    在Producer(发送端)，应用程序调用send()发送一条消息后，无论是同步发送还是异步发送，Kafka都不会立即把这条消息发送出去。它会先把这条消息存放在内存中缓存起来，然后选择合适的时机把缓存中所有的消息组成一批，一次性发送给Broker。通俗说就是攒一波再发送。
    在Broker(服务端)，Kafka不会把一批消息还原成多条消息，而是把每批消息当成一个“批消息”来处理。也就是说，在Broker的整个处理流程中，无论是写入磁盘、从磁盘中读出来、还是复制到其他副本，批消息都不会被解开，一直作为一个一条“批消息”来进行处理。
    在消费时，消息同样是以批为单位进行传递的，Consumer从Broker拉到一批消息后，在客户端把批消息解开，再一条一条交给应用程序代码处理。
    构件批消息和解开批消息分别在发送端和消费端的客户端完成，不仅减轻了Broker的压力，最重要的是减少了Broker处理请求的次数，提升了总体的处理能力。
二、使用顺序读写提升磁盘IO性能
    众所周知，对于磁盘来说，顺序读写的性能要远远好于随机读写。因为顺序读写只有一次寻址时间，而随机读写，每读一次都需要一次寻址时间。
    磁盘大致结构图如下：
    Kafka充分利用了磁盘的这个特性，它的存储设计的非常简单，对于每个分区，它把从Producer收到的消息，顺序的写入对应的log文件中，一个文件写满了就开启一个新的文件，这样顺序写下去。消费的时候，也是从某个全集的位置开始，也就是某一个log文件的某个位置，顺序的把消息读出来。
    这样一个简单的设计，充分利用了顺序读写的特性，大大的提升了Kafka在使用磁盘时的IO性能。
三、利用PageCache加速消息的读写
    通俗的说，PageCache就是操作系统在内存中给磁盘上的文件建立的缓存。无论我们使用何种语言编写的程序，在调用系统API读写文件时，并不会直接去读写磁盘上的文件，应用程序实际操作的都是PageCache，也就是文件在内存中的缓存的副本。
    应用程序在写入文件的时候，操作系统会先把数据写入到内存中的PageCache，然后在合适的时机再一批一批的写入到磁盘上。
    应用程序在读取文件时也是从PageCache中来读取数据，这时候就会出现两种情况。
    如果PageCache中有数据，那就直接读取，这样就免去了从磁盘上读取数据的时间；如果PageCache中没有数据，这时候操作系统就会触发一个缺页中断，应用程序的读取线程会被阻塞，操作系统将数据从磁盘文件中复制到PageCache中，然后应用程序再从PageCache中继续把数据读出来。
    应用程序在使用完某块PageCache后，操作系统不会立即将这个PageCache清除，而是尽可能的利用空闲的物理内存保存这些PageCache，除非系统内存不够用，操作系统才会清理掉一部分的PageCache，这就涉及到一些清理PageCache的策略，一般是使用LRU算法。
    Kafka就是利用了PageCache的这个特性，加速了消息读写速度。
四、ZeroCopy：零拷贝技术
    我们知道，在服务端，处理消费的大致逻辑是：
    1)从文件中找到消息数据，读到内存中；
    2)把消息通过网络发送给客户端。
    这个过程中，数据实际上做了2次或者3次的复制：
    1)从文件复制数据到PageCache中，如果命中PageCache，这一步可以省略；
    2)从PageCache复制到应用程序的内存空间；
    3)从应用程序的内存空间复制到Socket的缓冲区了，这个过程就是我们调用网络应用框架API发送数据的过程。
    Kafka使用零拷贝技术就可以把上面的复制次数减少一次，将上面的2、3步骤两次复制合并成一次复制，直接从PageCache中把数据复制到Socket缓冲区中，这样既减少一次数据复制，也不用CPU参与，直接由DMA控制器完成数据复制，速度更快。
======================================================
绿蚁新醅酒，红泥小火炉。
晚来天欲雪，能饮一杯无？
	

一、什么是字节？

  字节是存储数据的基本单位，并且还是硬件所能访问的最小单位。众所周知，CPU是不能直接访问硬盘上的数据，而只能直接地处理内存中的数据。硬盘上的数据必须先调入内存条中才可以运行。内存中存储数据的最小单位是“位”。字节是存储数据的基本单位。一个是基本单位，一个是最小单位，不要混淆了。

   内存里面存放的都是二进制代码。内存上有很多的“小格子”，每个“格子”中只能存放一个二进制数“0”或者“1”。一个“格子”就是一位，所以“位”要么是“0”或者“1”，不可能再有比“位”更小的单位了。


二、字节与位之间的关系？

  8个“小格子”也就是8“位”对应一个字节，即1Byte=8b。


三、为什么硬件所能访问的最小单位是字节，而不是位呢？

  因为硬件是通过地址总线访问内存的，而地址是以字节为单位分配的，所以地址总线只能精确到字节。那如何控制到字节的某一位，就要通过“位运算符”，即通过软件的方式来控制。


四、字节的换算？

  常见的存储单位主要有“bit（位）”、“B（字节）”、“KB（千字节）”、“MB（兆字节）”、“GB（千兆字节）”等。它们之间的换算关系如下：

  1、1字节(Byte) = 8位(bit)

  2、1KB( Kilobyte，千字节) = 1024B

  3、1MB( Megabyte，兆字节) = 1024KB

  4、1GB( Gigabyte，吉字节，千兆) = 1024MB

  5、1TB( Trillionbyte，万亿字节，太字节) = 1024GB

  6、1PB( Petabyte，千万亿字节，拍字节) = 1024TB

  7、1EB( Exabyte，百亿亿字节，艾字节) = 1024PB

  8、1 ZB(Zettabyte，十万亿亿字节，泽字节) = 1024EB

  9、1YB( Yottabyte，一亿亿亿字节，尧字节) = 1024ZB

  10、BB( Brontobyte，千亿亿亿字节)=1024YB

例如：计算机的内存是4GB，那么它能存放多少个0或者1（即能存放多少位）呢？

解答：由换算关系可知，1GB=1024MB=1024x1024KB=1024x1024x1024B=1024x1024x1024x8b

那么4GB也就是：4x1024x1024x1024x8（bit）

再由于一个1024=210，所以结果相当于32x230这么大！也就是一个4GB内存条所能存储的数据。

类似，对于一个500GB的硬盘，它所能存储的数据就是500x1024x1024x1024x8位，即能存放4000个230这么多个“0”或者“1”。


五、小常识解惑？

  比如在电脑城里买了一个带有500GB硬盘的笔记本，但是真正显示出来的硬盘大小却没有500GB，而是略小于500GB的，为什么呢？

  解惑：原因在于计量单位不一样！在商场买的硬盘或者U盘都是以1000为单位的，而计算机是以1024为单位的，所以显示的就没有那么大。


大家可以自己计算一下：将16GB的空U盘插入电脑后，显示的大小略小于16GB，那么它的实际大小大约是多少呢？


六、小扩展

两个小概念了解一下：“字”和“字长”

  字：在计算机中，一串数码作为一个整体来处理或运算的，称为一个计算机字，简称字。字通常分为若干个字节（每个字节一般是8位）。在存储器中，通常每个单元存储一个字。因此每个字都是可以寻址的。字的长度用位数来表示 。

  字长：计算机的每个字所包含的位数称为字长，计算的字长是指它一次可处理的二进制数字的数目。一般地，大型计算机的字长为32-64位，小型计算机为12-32位，而微型计算机为4-16位。字长是衡量计算机性能的一个重要因素 。

1.进程和线程的基本概念
进程:
进程可以理解为一个程序的实例,就可以理解为一个应用或程序,
当程序被运行时,会把磁盘中的程序代码加载到内存，同时开启一个进行;
线程：
可以理解为进程中的一条条指令流，线程是进程的子集,将一条条的指令按照一定的顺序交给CPU,
线程是Java中最小的调度单位。

2.并行与并发的概念
并发：同一个CPU内核处理不同的指令,并发（concurrent）是同一时间应对（dealing with）多件事情的能力;
并行：不同的CPU内核同时处理不同的指令,并行（parallel）是同一时间动手做（doing）多件事情的能力;




3.创建线程的三种方式：
     /**
     * 方法一：
     * 直接使用Thread匿名内部类的形式
     * 创建线程，调用线程Thread类的start方法执行线程体
     */
    @Test
    public void createThread1() {
        Thread th = new Thread(() -> {
            log.debug("线程体开始执行");
            try {
                Thread.sleep(1);
            } catch (InterruptedException e) {
                e.printStackTrace();
            }
            log.debug("线程方法执行结束");
        });
        th.start();
    }

     /**
     * 方法二：
     * 使用Runnable接口，
     * 使用匿名内部类，重写run方法
     * 把创建的runnable接口引用放入Thread示例
     * 调用Thread的start方法执行线程体
     */

    public void createThread2() {
        Runnable r = () -> {
            log.info("to do someThing");
        };
        Thread re = new Thread(r);
        re.start();
    }


     /**
     * 方式三：
     * FutureTask 配合 Thread
     */
    @Test
    public void createThread3() throws ExecutionException, InterruptedException {
        // 创建任务对象
        FutureTask<Integer> task3 = new FutureTask<>(() -> {
            log.debug("hello");
            return 100;
        });
        // 参数1 是任务对象; 参数2 是线程名字，推荐
        new Thread(task3, "t3").start();
        // 主线程阻塞，同步等待 task 执行完毕的结果
        Integer result = task3.get();
        log.debug("结果是:{}", result);
    }

4.查看进程的方法
windows

任务管理器可以查看进程和线程数，也可以用来杀死进程

tasklist 查看进程

taskkill 杀死进程

linux

ps -fe 查看所有进程

ps -fT -p  查看某个进程（PID）的所有线程

kill 杀死进程

top 按大写 H 切换是否显示线程

top -H -p  查看某个进程（PID）的所有线程