纯新手，问题在大家看来可能很简单，因自学，希望大家帮帮忙，别嫌弃。:loveliness:问题是这样的：
matrixMulCUDA<<<grid, threads >>>(d_C, d_A, d_B, dimsA.x, dimsB.x);
。。。
matrixMulCUDA（）
{
。。。。。
}
想问，每一次使用matrixMulCUDA 时，参与计算的线程应该只有一个吧？如果是这样那么在矩阵乘法例子里有这样的代码：shared float As[BLOCK_SIZE][BLOCK_SIZE]；
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
As[ty][tx] = A[a + wA * ty + tx];
Bs[ty][tx] = B[b + wB * ty + tx];
说是可以把一个子矩阵存入共享内存，可是如果只有一个线程在工作，没有循环是怎么存进去的呢？
参与计算的线程数是由你的grid,threads决定的，如果是1,1那么只有一个线程。你看一下SDK的例子，里面的确有循环再读取矩阵的数据！
LZ您好，不知道您“每一次使用matrixMulCUDA 时，参与计算的线程应该只有一个吧？”这种说法的来源是？
我来大致说一下。
您每次在跑kernel的时候，总的线程数是 “每个block内的线程数” * “grid内的block数量”。
分别对应您的代码中的“threads”和“grid”两个量。
这两个量是dim3类型的，最多可以有3个分量。那么您的总的乘积是两者自身3个分量的乘积再相乘。（对于您没有指定的分量，默认值是1；以及如果您使用一维，那么可以用int类型的变量代替）
shared memory是每个block独有的，那么每个block内可以用全部或者部分线程来完成shared memory数据的赋值工作。
这取决于您的程序实现和算法安排。
不过一般地，不使用<<<1,1>>>这种安排，这样效率非常低。
所以，您可以根据您的具体代码，考虑一下shared memory是怎么赋值的。
如果没有循环/展开的循环，而完成了一组赋值的话，那么应当是多个线程在干活的。
最后，我猜测一下您的“一个线程”说法的来源，您的“一个线程”可能表示的是“一个主机端线程”。但是在CUDA里面，一般说的线程，都是GPU上跑的device端的线程，一般是数量庞大的，也是来干活的。
希望上述解释可以解决LZ的疑惑，祝LZ好运！
:)太感谢了，我明白了，对这个东西很感兴趣，但是刚开始很多地方理解的不到位。
matrixMulCUDA(float *C, float *A, float *B, int wA, int wB)
{
int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;
int aBegin = wA * BLOCK_SIZE * by;
int aEnd = aBegin + wA - 1;
int aStep = BLOCK_SIZE;
int bBegin = BLOCK_SIZE * bx;
int bStep = BLOCK_SIZE * wB;
float Csub = 0;
for (int a = aBegin, b = bBegin;a <= aEnd;a += aStep, b += bStep)
{
shared float As[BLOCK_SIZE][BLOCK_SIZE]；
shared float Bs[BLOCK_SIZE][BLOCK_SIZE];
As[ty][tx] = A[a + wA * ty + tx];
Bs[ty][tx] = B[b + wB * ty + tx];
__syncthreads();
#pragma unroll
for (int k = 0; k < BLOCK_SIZE; ++k)
{
Csub += As[ty][k] * Bs[k][tx];
}
__syncthreads();
C[c + wB * ty + tx] = Csub;
}
具体的代码是这样的，这属于版主说的多个线程在工作么
matrixMulCUDA(float *C, float *A, float *B, int wA, int wB)
{
int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;
int aBegin = wA * BLOCK_SIZE * by;
int aEnd = aBegin + wA - 1;
int aStep = BLOCK_SIZE;
int bBegin = BLOCK_SIZE * bx;
int bStep = BLOCK_SIZE * wB;
float Csub = 0;
for (int a = aBegin, b = bBegin;a <= aEnd;a += aStep, b += bStep)
{
shared float As[BLOCK_SIZE][BLOCK_SIZE]；
shared float Bs[BLOCK_SIZE][BLOCK_SIZE];
As[ty][tx] = A[a + wA * ty + tx];
Bs[ty][tx] = B[b + wB * ty + tx];
__syncthreads();
#pragma unroll
for (int k = 0; k < BLOCK_SIZE; ++k)
{
Csub += As[ty][k] * Bs[k][tx];
}
__syncthreads();
C[c + wB * ty + tx] = Csub;
}
具体的代码是这样的，这属于版主说的多个线程在工作么
LZ您好，如果您不较真说发布kernel的时候用的是<<<1,1>>>的话，这个代码显然是（device端的）多线程的。
代码中有标记block内线程编号的tx和ty，还有标记block编号的bx和by。
代码中无论是shared memory赋值，还是相乘计算都是多线程的。
请LZ明察。
恩:)还有个地方不太明白，麻烦版主。程序里调用内核计算时，有这样第一段代码
int nIter = 300;
for (int j = 0; j < nIter; j++)
{
if (block_size == 16)
{
matrixMulCUDA<16><<< grid, threads >>>(d_C, d_A, d_B, dimsA.x, dimsB.x);
}
else
{
matrixMulCUDA<32><<< grid, threads >>>(d_C, d_A, d_B, dimsA.x, dimsB.x);
}
之前程序里还定义了， 矩阵A和B尺寸
dim3 dimsA(52block_size, 52block_size, 1);
dim3 dimsB(54block_size, 52block_size, 1);
线程块包含的线程数 dim3 threads(block_size, block_size)；
以及线程块数 dim3 grid(dimsB.x / threads.x, dimsA.y / threads.y);
那么就应该是分出了545*2个线程块，那么最开始为什么是 int nIter = 300？是我理解出现了偏差么。麻烦你了。
在指南里，我看到关于这个语句的解释有些胡度，比如 x<<<5, 16>>>(A, B, C);这条语句出现的时候是一块线程 16个先一起执行x，然后下一块再执行，还是这条语句一出现，就是5块线程一起执行x.
LZ您好，这里int nlter=300是给下面的循环用的，而下面的循环是将kernel用同样的参数启动了300次。
所以这个300和您kernel里面启动了多少个线程是没有关系的。
那么，为何要使用这个循环呢？这不是无用功么？
是的，我估计这个循环是用来将原kernel原样跑300次，然后平均下下求每次计算时间的。因为这个示例的kernel可能规模太小，不好计时。
循环里面的if-else判断是怎么回事？
原程序应该是写了两个不同的kernel，一个是为了1616的block写的，一个是3232的block写的。
在早期的GPU上，一个block最大是512个线程，放不下32*32的block。
所以前面可能根据检测GPU版本的结果，选择了不同的block大小，并加以计算。
大致如上，供您参考，祝您好运~
LZ您好，您的问题是在问发射kernel的时候，那些线程和线程块是同时拿去执行的还是先后执行的。
说明这个问题，需要先简单回顾一下GPU的硬件组成，以及线程在GPU上执行时的安排。
我们知道，线程是按照warp为单位在GPU的SM（SMX）上执行的。一个GPU可以有一个或者多个SM，这视GPU的规模不同而不同，同架构，越高档的GPU ，SM数量越多，不同架构的GPU不直接比较。
一个SM（SMX）上可以resident一个或者多个block，具体每个SM上能resident多少block，能resident多少thread，一个block上最多多少thread，等限值可以参看 CUDA C Programming Guide的Appendix.F
如果block很小（block内thread数量小），那么SM上能resident的block数量可能先达到饱和。
如果block很大或者使用shared memory等资源过多，那么SM上也无法加载更多的block。
在程序具体的资源需求和GPU具体架构的硬件下，一个SM上能resident的block数量是固定的。
此时，如果您的block总数大于您GPU上的所有SM所能resident的block数量，那么后面的block需要等前面的block结束以后，才能加载到GPU上继续进行计算。实际应用中，block数量和thread数量一般是巨大的，会铺满整个GPU，并前赴后继地运行。
如果您的block总数小于您GPU上的所能resident的block总数，那么一开始就都会被发布到SM上执行，不会有等待的。
同时需要说明的是，前述的情况对于程序员而言，某种程度上说是透明的。您可以不管您的GPU到底有几个SM。如果是一代神U，那么可能立即全部发布了，如果战斗力只有5的渣U，可能需要等前面的block计算完，后面的才能上。但是这仅仅是慢点，逻辑上并没有差别。
------------------------------------纠结的分割线-------------------------------------------------
如果您需要纠结一个block内部在SM上是如何计算的，有没有先后顺序，那么我再说几句。
SM内部一次仅能对少数的warp的当前1~2条命令进行计算，少数目前可能指1,2,4。
当前这些warp的这些指令计算完（灌入流水线完）之后，可能继续计算不相关的指令，也可能被切换出去换其他就绪的warp进来计算。GPU就是依靠这种切换来使用计算掩盖延迟的。同时这种切换是没有顺序保证的，不保证谁一定在谁前面。
所以，从这个更细节的角度考虑，一个SM上resident的block，其实是按照warp为单位，乱序进SM计算的。既不是顺序，也不是完全同时。
-------------------------------------总结的分割线--------------------------------------
综上所述，LZ可以考虑下您的具体情况是怎么回事。
您在9#中给出的示例，block和thread规模都很小，一般来说只用于示例。
以及您没有给出您kernel的具体实现，所以无法给您确切的结论。
不过根据上述讨论，您可以自行判断。
欢迎您莅临CUDA ZONE，祝您编码顺利~
恩恩，是这样啊，我一直在想怎么重复在做呢，这下明白了，:)谢谢版主
恩恩明白了，因为他们是无序的，所以在程序中会利用同步设置是他们达到同步的效果。我看到指南里说，一个块能包含的最大线程数就是1024 ，是不是说当块内的线程数达到1024后，想增大规模就只能增加块数了？还有就这个例子而言，分块减少了线程在显存里读取的次数，提高了效率，如果我想在此基础上，进一步提高的话，我有两个设想，不知道是否合理：一、例子里面是一个线程执行内核时，用双重循环的方式求出矩阵C的一个元素的值，可不可以子块内的线程把子块的数据保存到共享内存后，采取对加的方式，也就是一行的第一个线程读取元素加上最后一个读取的，这样继续下去是不是可以减少运算次数提高效率。二、能不能在内核中直接把主存里保存的A和B矩阵的数据读入共享内存，我记得指南上说，数据是可以存放在显存或者共享内存里的，不知道我这样想是否可行:)
LZ您好，我觉得您的理解尚有偏差。
1：__syncthreads()只能同步一个block内部的线程，block之间依然是不保证顺序的。
2：一个block最大线程数是1024没错（对于目前的硬件而言如此，老硬件更少），大规模的计算需要更多的block甚至更多的grid（多启动kernel）。但并非我们需要把block塞到1024线程，这并非最佳选择。block大小选定和很多因素有关，一般典型的大小是192到256线程。以及您可以在线程数减少的时候，增加单个线程的工作量。
3：您引用的例子只是用来示例cuda编程的，程序的 效能取决于很多原因，示例中也许并非是一个高效的框架，您可以尝试改写和比较，但是一般来说，如果您需要矩阵乘法，那么现成的库的速度会比您自己写的快。
4：您说的第一种优化方案，我没看明白，无法评价，您可以自己测试一下速度。另外，每个线程读取到自己的寄存器里面的数据是私有的，无法直接跨线程相互访问。可以通过shared memory交互。
5：这个做法是不可行的，因为shared memory容量非常小，而需要计算的矩阵往往较大。而且，shared memory的有效范围是一个block，保存该block不处理的数据也没有意义。shared memory 的大小是每SM 16KB（对于1.x硬件）或每SM 16KB/48KB可调（对于2.x硬件），在kepler中增加了32KB选项。总之最大值也不过每SM 48KB，是作为手工可控的高速cache来用的。2.x和3.x硬件每SM/SMX提供的L1 cache和shared memory的总数是64KB，设置为shared memory 16KB 的时候，L1cache 为64-16=48KB，反之亦然。
祝您编码愉快~
我的想法没说清楚，是这样的，我是想把当前这块线程要处理的数据从主存里直接读入这块线程的共享内存。 也就是像这个矩阵相乘的例子里，是把矩阵A和B按照blocksize*blocksize的大小分成了多个子矩阵，相应的分出同样数量的线程块，例子里是先在主存里分配A和B的空间，赋值，然后拷贝到显存中，最后由线程块，读入共享内存，加以计算，我是想能不能，把每个线程块对应计算的子矩阵直接从主存读入共享内存，这样省不是去了拷贝到显存中，在读出的过程，不知道能不能相对节省时间。
LZ您好，如果您确定要这样，可以使用主机端的页锁定内存空间，并使用zerocopy技术来实现。
但是这样未必会提升速度。
如果需要掩盖传输造成的时间，那么可以将一个大的数组分为多个小一些的数组，并使用stream的方式，实现计算和copy互相掩盖，并且编程实现要更为复杂一些。
祝您编码顺利~
版主的意思就是 把数据从主存拷贝到显存时，原本是把矩阵A和B各保存在一个大数组里，然后传输，可以把他们拆分成一些小数组，在传递？
您查看一些使用stream 来掩盖传输与计算的例子即可。
恩恩:)谢谢版主
嗯嗯，不客气的，欢迎您常来论坛~