初学，对照指南有一些关于GPU硬件实现和内存结构的理解，希望版主和大家帮忙指点。
1、指南里讲到我们可以把GPU描述为一组多处理器。是不是说我们把GPU中的多核分成多组，每一组称之为一个多处理器。
2、指南里提到每个多处理器 拥有的片上内存 包括：多组32位的寄存器（一个处理器有一组32位的寄存器），一块并行数据缓存（共享内存），一块常量高速缓存和一块纹理高速缓存。是不是可以这样认为，GPU中的多处理器平分了片上的共享内存和常量缓存还有纹理缓存。
3、指南中提到在处理线程块时，把他们成批次的发给每个多处理器，多处理器接收到一个线程块就要从自己的共享内存空间里分出一分来给这块线程，同时还要分出寄存器给块内所有线程，是不是说一个多处理器一批次能接受多少线程块，要看他的共享内存空间够不够分。
4、指南里讲到设备的DRAM分为 local memory global memory constant memory texture memory.是不是可以这样理解 local memory 是直接面向每个线程，需要的话就在他上面分出一块给线程，而 global memory constant memory texture memory是面向每个网格的 也就是产生一个线程网格 就要分别在他们上面分出一块属于这个线程网格的空间。
反复阅读指南的理解 望版主和大家帮助。
LZ您好，您的问题大致回答如下：
1：是的，基本如此，但一般是反过来叙述，一组SP/CUDA Core和其他一些逻辑单元组成一个SM，多个SM和一些全局功能的单元组成整个GPU芯片。
2：是的，但是和1：的情况一样，一般是按照SM堆积为GPU的角度理解，而不是反过来考虑SM瓜分GPU的资源。
3：SM上能resident多少个blocks有多个制约的因素，主要有，block的线程数多少，SM上最多resident block的数量，寄存器的使用量，shared memory的使用量等。您说的shared memory限制只是其中的一个约束。（同时，如果kernel不使用shared memory的话，则不分配shared memory给block）
4：上述4种memory都是使用显存的DRAM的，但是其中local memory是每个线程私有的，其他几种是grid的所有线程都可以访问的。至于划分出空间，还需要您通过API申请使用，才会分配空间。比如您使用cudaMalloc()申请一段global memory，此时才会分配一段空间给您使用，否则DRAM的空间是闲置备用的。
大致如上，供您参考。
祝您编码愉快~
太感谢了:)
不客气的，欢迎您常来论坛转转~
祝您编码愉快~
我试着改写了 新建工程里默认的那个向量相加的程序，想在内核里引进共享内存结果出现错误，改写的代码如下
include “cuda_runtime.h”
include “device_launch_parameters.h”
include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
template global void
addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
shared int m [arraySize];
shared int n [arraySize];
m[i]= a[i];
n[i]= b[i];
__syncthreads();
c[i] = m[i] +n[i];
}
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };
…
说是参数传递的问题，我是仿照矩阵乘法的例子改写的，不知道错在哪里
LZ您好，请您用“代码模式”重发您的代码，直接贴到正文里面会导致很多字符丢失的问题。
同时，请补充addWithCuda()函数的定义。
以及，请描述一下您的问题是？结果不正确？程序崩溃？还是神马？
祝您debug顺利~
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
template <int arraySize>__global__ void
addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
__shared__ int m[arraySize];
__shared__ int n[arraySize];
m[i]= a[i];
n[i]= b[i];
__syncthreads();
c[i] = m[i] + n[i];
}
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };
// Add vectors in parallel.
cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCuda failed!");
return 1;
}
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
c[0], c[1], c[2], c[3], c[4]);
// cudaDeviceReset must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceReset failed!");
return 1;
}
return 0;
}
这是addWithCuda()函数代码
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
// Launch a kernel on the GPU with one thread for each element.
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
// cudaDeviceSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
楼主，你kernel中的arraysize在哪呢？都没定义的。
我是仿照矩阵乘法的这段代码改的，其实用别的方法很容易就改成功了，可是我就是不太明白为什么矩阵这个代码里这样写可以，
template <int BLOCK_SIZE> __global__ void
matrixMulCUDA(float *C, float *A, float *B, int wA, int wB)
{
// Block index
int bx = blockIdx.x;
int by = blockIdx.y;
// Thread index
int tx = threadIdx.x;
int ty = threadIdx.y;
// Index of the first sub-matrix of A processed by the block
int aBegin = wA * BLOCK_SIZE * by;
// Index of the last sub-matrix of A processed by the block
int aEnd = aBegin + wA - 1;
// Step size used to iterate through the sub-matrices of A
int aStep = BLOCK_SIZE;
// Index of the first sub-matrix of B processed by the block
int bBegin = BLOCK_SIZE * bx;
// Step size used to iterate through the sub-matrices of B
int bStep = BLOCK_SIZE * wB;
// Csub is used to store the element of the block sub-matrix
// that is computed by the thread
float Csub = 0;
// Loop over all the sub-matrices of A and B
// required to compute the block sub-matrix
for (int a = aBegin, b = bBegin;
a <= aEnd;
a += aStep, b += bStep)
{
// Declaration of the shared memory array As used to
// store the sub-matrix of A
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
// Declaration of the shared memory array Bs used to
// store the sub-matrix of B
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
// Load the matrices from device memory
// to shared memory; each thread loads
// one element of each matrix
As[ty][tx] = A[a + wA * ty + tx];
Bs[ty][tx] = B[b + wB * ty + tx];
// Synchronize to make sure the matrices are loaded
__syncthreads();
// Multiply the two matrices together;
// each thread computes one element
// of the block sub-matrix
#pragma unroll
for (int k = 0; k < BLOCK_SIZE; ++k)
{
Csub += As[ty][k] * Bs[k][tx];
}
// Synchronize to make sure that the preceding
// computation is done before loading two new
// sub-matrices of A and B in the next iteration
__syncthreads();
}
// Write the block sub-matrix to device memory;
// each thread writes one element
int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;
C[c + wB * ty + tx] = Csub;
}
在矩阵乘法的内核函数里也没有对 BLOCK_SIZE进行定义 但是可以用 我在我的程序里这样写就不行
LZ您好，您既然是参考的matrixMul的，那么原程序中调用相乘函数的时候的用法也请您参考一下。
在CUDA 5的例子里面，VS 2008版本，250行附近。
估计您一看便知。
祝您debug顺利~
对C++不太熟悉，以我有限的C++知识，LZ你对模板的使用有点问题。
好好看看模板的使用。
看了那段代码，可是在内核函数里面确实没有对BLOCK_SIZE进行定义就直接使用了，所以不太明白为什么。
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
改为
addKernel<5><<<1, size>>>(dev_c, dev_a, dev_b);
即可。
您可以回头再看看matrixMul的代码。
祝您修改顺利~