6
在開普勒硬件上分析Visual Profiler中的內核時，我注意到Profiler顯示全局加載和存儲緩存在L1中。 我很困惑，因爲節目指南和開普勒調諧該手動狀態：本地存儲器訪問中的GPU開普勒Kepler全局內存訪問和L1緩存
L1高速緩存只保留，諸如寄存器 溢漏和堆棧數據。全局負載僅緩存在L2中（或緩存中的只讀數據 ）。
沒有寄存器溢出（profiler顯示L1緩存，即使是原始的2行'add'內核），我不確定'堆棧數據'在這裏意味着什麼。
GK110白皮書顯示，除了一種情況外，全局訪問都將通過L1緩存：通過只讀緩存（__ldg）加載。 這是否意味着雖然全局訪問通過L1硬件，但它們實際上並沒有被緩存？這是否也意味着如果我將溢出的寄存器數據緩存在L1中，這些數據可能會因gmem訪問而被驅逐？
更新：我已經意識到我可能會誤讀分析器是給我的信息，所以這裏是內核代碼以及探查的結果（我用相同的嘗試都在Titan和K40結果）。
template<typename T>
__global__ void addKernel(T *c, const T *a, const T *b)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
c[i] = a[i] + b[i];
}
...
// Kernel call
float* x;
float* y;
float* d;
// ...
addKernel<<<1024, 1024>>>(d, x, y);
cudaError_t cudaStatus = cudaDeviceSynchronize();
assert(cudaSuccess == cudaStatus);
視覺探查輸出：
L1號碼來撥打給L1高速緩存完美感啓用GMEM訪問。對於負載，我們有：
65536 * 128 = = 2 * 4 * 1024 * 1024
更新2：添加SASS和PTX代碼。 SASS代碼非常簡單，包含從常量內存讀取以及從/到全局內存（LD/ST指令）的加載/存儲。
Function : _Z9addKernelIfEvPT_PKS0_S3_
.headerflags @"EF_CUDA_SM35 EF_CUDA_PTX_SM(EF_CUDA_SM35)"
/* 0x088cb0a0a08c1000 */
/*0008*/ MOV R1, c[0x0][0x44]; /* 0x64c03c00089c0006 */
/*0010*/ S2R R0, SR_CTAID.X; /* 0x86400000129c0002 */
/*0018*/ MOV32I R5, 0x4; /* 0x74000000021fc016 */
/*0020*/ S2R R3, SR_TID.X; /* 0x86400000109c000e */
/*0028*/ IMAD R2, R0, c[0x0][0x28], R3; /* 0x51080c00051c000a */
/*0030*/ IMAD R6.CC, R2, R5, c[0x0][0x148]; /* 0x910c1400291c081a */
/*0038*/ IMAD.HI.X R7, R2, R5, c[0x0][0x14c]; /* 0x93181400299c081e */
/* 0x08a0a4b0809c80b0 */
/*0048*/ IMAD R8.CC, R2, R5, c[0x0][0x150]; /* 0x910c14002a1c0822 */
/*0050*/ IMAD.HI.X R9, R2, R5, c[0x0][0x154]; /* 0x931814002a9c0826 */
/*0058*/ LD.E R3, [R6]; /* 0xc4800000001c180c */
/*0060*/ LD.E R0, [R8]; /* 0xc4800000001c2000 */
/*0068*/ IMAD R4.CC, R2, R5, c[0x0][0x140]; /* 0x910c1400281c0812 */
/*0070*/ IMAD.HI.X R5, R2, R5, c[0x0][0x144]; /* 0x93181400289c0816 */
/*0078*/ FADD R0, R3, R0; /* 0xe2c00000001c0c02 */
/* 0x080000000000b810 */
/*0088*/ ST.E [R4], R0; /* 0xe4800000001c1000 */
/*0090*/ EXIT ; /* 0x18000000001c003c */
/*0098*/ BRA 0x98; /* 0x12007ffffc1c003c */
/*00a0*/ NOP; /* 0x85800000001c3c02 */
/*00a8*/ NOP; /* 0x85800000001c3c02 */
/*00b0*/ NOP; /* 0x85800000001c3c02 */
/*00b8*/ NOP; /* 0x85800000001c3c02 */
PTX：
.visible .entry _Z9addKernelIfEvPT_PKS0_S3_(
.param .u64 _Z9addKernelIfEvPT_PKS0_S3__param_0,
.param .u64 _Z9addKernelIfEvPT_PKS0_S3__param_1,
.param .u64 _Z9addKernelIfEvPT_PKS0_S3__param_2
)
{
.reg .s32 %r<5>;
.reg .f32 %f<4>;
.reg .s64 %rd<11>;
ld.param.u64 %rd1, [_Z9addKernelIfEvPT_PKS0_S3__param_0];
ld.param.u64 %rd2, [_Z9addKernelIfEvPT_PKS0_S3__param_1];
ld.param.u64 %rd3, [_Z9addKernelIfEvPT_PKS0_S3__param_2];
cvta.to.global.u64 %rd4, %rd1;
.loc 1 22 1
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r1, %r2, %r3;
cvta.to.global.u64 %rd5, %rd2;
mul.wide.s32 %rd6, %r4, 4;
add.s64 %rd7, %rd5, %rd6;
cvta.to.global.u64 %rd8, %rd3;
add.s64 %rd9, %rd8, %rd6;
.loc 1 23 1
ld.global.f32 %f1, [%rd9];
ld.global.f32 %f2, [%rd7];
add.f32 %f3, %f2, %f1;
add.s64 %rd10, %rd4, %rd6;
.loc 1 23 1
st.global.f32 [%rd10], %f3;
.loc 1 24 2
ret;
}
看看[這篇文章]（http://stackoverflow.com/questions/19627702/l2-cache-in-kepler）。 L1緩存和只讀緩存是不同的。在評論中，羅伯特曾提到「開普勒通常不會爲普通的全球負載啓用L1」。 – Farzad
感謝您的鏈接。我明白L1和紋理緩存是不同的。除了編程指南中提到的內容（即寄存器溢出和堆棧數據）之外，我對Robert的「正常」意義很感興趣。 –
你可以顯示你的.PTX輸出嗎？我在想你的短內核可能會有低級指令，它們利用L1高速緩存來保存中間結果。 – Farzad