我有以下简单代码：
全局内存负载效率的价值让我感到惊讶。由于我相信我正在执行完全联合的内存访问，所以我本来期望在两种情况下都有100％的效率。那么我的问题是：
当我执行联合内存访问时，为什么我有50％的全局内存负载效率，而我却有100％的全局内存存储效率？
我还调查了其他可能有用的指标/事件，这些指标/事件也可以报告。
上述两个事件似乎相互矛盾，根据我的（可能是错误的）理解。在
这似乎与我正在执行的完全合并的内存访问事实一致。
反汇编代码如下：
现在Visual Profiler报告了一个未缓存的全局加载事务，但没有
编辑
我对此问题进行了更深入的调查，增加了
当块数为奇数时，即我的GT540M卡的两个流多处理器的负载不平衡时，全局内存加载效率小于
如果我通过使用
一些图片展示了全局内存加载效率的不同值的行为。
通过禁用
#include<stdio.h>
#define BLOCKSIZE_X 32
#define BLOCKSIZE_Y 1
int iDivUp(int a, int b) { return ((a % b) != 0) ? (a / b + 1) : (a / b); }
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
__global__ void kernel0(float *d_a, float *d_b, const unsigned int M, const unsigned int N)
{
const int tidx = threadIdx.x + blockIdx.x * blockDim.x;
const int tidy = threadIdx.y + blockIdx.y * blockDim.y;
if ((tidx < M)&&(tidy < N)) {
d_b[tidy * M + tidx] = d_a[tidy * M + tidx];
}
}
void main()
{
const unsigned int M = 32;
const unsigned int N = 1;
float *d_a; cudaMalloc((void**)&d_a, M*N*sizeof(float));
float *d_b; cudaMalloc((void**)&d_b, M*N*sizeof(float));
dim3 dimGrid(iDivUp(M, BLOCKSIZE_X), iDivUp(N, BLOCKSIZE_Y));
dim3 dimBlock(BLOCKSIZE_X, BLOCKSIZE_Y);
kernel0<<<dimGrid, dimBlock>>>(d_a, d_b, M, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
cudaDeviceReset();
}
这段代码执行两个包含32个浮点数的数组之间的赋值操作。我正在尝试理解全局内存协同访问与全局内存读/写效率以及其他指标/事件之间的关系。
可视化分析器显示以下指标：
Global Memory Load Efficiency = 50%
Global Memory Store Efficiency = 100%
全局内存负载效率的价值让我感到惊讶。由于我相信我正在执行完全联合的内存访问，所以我本来期望在两种情况下都有100％的效率。那么我的问题是：
当我执行联合内存访问时，为什么我有50％的全局内存负载效率，而我却有100％的全局内存存储效率？
我还调查了其他可能有用的指标/事件，这些指标/事件也可以报告。
gld_inst_32bit = 32 (Number of 32-bit global memory load transactions)
gst_inst_32bit = 32 (Number of 32-bit global memory store transactions)
实际上我正在请求加载/写入 32位浮点数。
uncached global load transaction = 0 (Number of uncached global load transactions)
l1 global load miss = 2 (Number of global load misses in L1 cache)
上述两个事件似乎相互矛盾，根据我的（可能是错误的）理解。在
l1 缓存未命中的情况下，我预期第一个事件与 0 不同。gld_request = 1 (Number of executed global load instructions per warp in a SM)
gst_request = 1 (Number of executed global store instructions per warp in a SM)
这似乎与我正在执行的完全合并的内存访问事实一致。
反汇编代码如下：
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R3, SR_CTAID.Y; /* 0x2c0000009800dc04 */
/*0010*/ S2R R4, SR_TID.Y; /* 0x2c00000088011c04 */
/*0018*/ IMAD R4, R3, c[0x0][0xc], R4; /* 0x2008400030311ca3 */
/*0020*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0028*/ ISETP.LT.U32.AND P0, PT, R4, c[0x0][0x2c], PT; /* 0x188e4000b041dc03 */
/*0030*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0038*/ IMAD R0, R0, c[0x0][0x8], R2; /* 0x2004400020001ca3 */
/*0040*/ ISETP.LT.U32.AND P0, PT, R0, c[0x0][0x28], P0; /* 0x18804000a001dc03 */
/*0048*/ @!P0 BRA.U 0x78; /* 0x40000000a000a1e7 */
/*0050*/ @P0 IMAD R2, R4, c[0x0][0x28], R0; /* 0x20004000a04080a3 */
/*0058*/ @P0 ISCADD R0, R2, c[0x0][0x20], 0x2; /* 0x4000400080200043 */
/*0060*/ @P0 ISCADD R2, R2, c[0x0][0x24], 0x2; /* 0x4000400090208043 */
/*0068*/ @P0 LD R0, [R0]; /* 0x8000000000000085 */
/*0070*/ @P0 ST [R2], R0; /* 0x9000000000200085 */
/*0078*/ EXIT; /* 0x8000000000001de7 */
编辑
我的配置：CUDA 6.5，GeForce GT540M，Windows 7。
如果我将M从32增加到64，启动两个块并使我的卡的两个可用流多处理器忙碌，那么全局内存负载效率就会变为100％，这些是新的指标/事件：
gld_inst_32bit = 64
gst_inst_32bit = 64
uncached global load transaction = 0
l1 global load miss = 2
gld_request = 2
gst_request = 2
预计增加gld_inst_32bit、gst_inst_32bit、gld_request和gst_request是合理的，因为现在我正在加载/存储64个float和2个全局内存加载/存储协同请求。但我仍然不明白为什么uncached global load transaction和l1 global load miss保持不变，而全局内存加载吞吐量发生变化以提供100%的效率。
编辑
M=32在Kepler K20c上的结果：
Global Memory Load Efficiency = 100%
Global Memory Store Efficiency = 100%
gld_inst_32bit = 64
gst_inst_32bit = 64
gld_request = 1
gst_request = 1
uncached global load transaction = 1
l1 global load miss = 0
l1 global load hit = 0
现在Visual Profiler报告了一个未缓存的全局加载事务，但没有
l1全局加载缺失。编辑
我对此问题进行了更深入的调查，增加了
M的值并保持BLOCKSIZE_X不变。当块数为奇数时，即我的GT540M卡的两个流多处理器的负载不平衡时，全局内存加载效率小于
100％，否则在偶数情况下为100％。全局内存加载效率在奇数情况下随着块数的增加而缓慢趋向于100％。如果我通过使用
-Xptxas -dlcm=cg来禁用L1缓存，如@Jez建议的那样编译，则全局内存加载效率始终等于100％，因为它是全局内存存储效率。我知道全局内存存储不使用L1缓存，而只使用L2。一些图片展示了全局内存加载效率的不同值的行为。
M=32
M=64
M=96
M=128
M=160
M=192
M是32的整数倍，以便单个warp可以加载整个缓存行。通过禁用
L1，我已经使得：M=32。
M=64
M=96
编辑 - 特斯拉 C2050 的结果
M = 32 33.3%
M = 64 28.6%
M = 96 42.9%
M = 128 57.1%
M = 160 71.4%
M = 192 85.7%
M = 224 100%
M = 256 114%
M = 288 90%
如果我禁用 L1 缓存，所有情况下的全局内存加载效率都达到了 100%。
-Xptxas -dlcm=cg进行编译，则会生成未在L1中缓存的加载，并且“未缓存的全局加载事务”计数器将增加。注意：这仅涉及L1，而不涉及L2。 - Jezd_a）最初不在 Fermi 情况下 SM 的 L1 缓存中，因此出现了 L1 缓存未命中。这是一个“未命中”。在 Kepler 情况下，全局加载禁用了 L1 缓存，因此所有加载都是“未缓存的”，并且没有 L1 参与，因此不会出现 L1 未命中。 - Robert Crovella-Xptxas -dlcm=cgзј–иҜ‘пјҢзҰҒз”ЁL1зј“еӯҳпјҲжӯЈеҰӮ@Jezе»әи®®зҡ„йӮЈж ·пјүпјҢеҲҷе…ЁеұҖеҶ…еӯҳеҠ иҪҪж•ҲзҺҮе§Ӣз»ҲзӯүдәҺ100пј…пјҢе°ұеғҸе…ЁеұҖеҶ…еӯҳеӯҳеӮЁж•ҲзҺҮдёҖж ·гҖӮжҲ‘зҹҘйҒ“е…ЁеұҖеҶ…еӯҳеӯҳеӮЁдёҚдҪҝз”ЁL1зј“еӯҳпјҢд»…дҪҝз”ЁL2зј“еӯҳпјҢжӯЈеҰӮжӮЁеҗ‘жҲ‘жҢҮеҮәзҡ„йӮЈж ·Keplerзҡ„е…ЁеұҖеҶ…еӯҳеҠ иҪҪгҖӮзҗҶи§ЈиҝҷдёҖзӮ№зҡ„е…ій”®дјјд№ҺжҳҜL1 / L2зј“еӯҳпјҹ - Vitality