如何将所有访问指向CUDA中的全局内存？

我希望我的程序能够访问全局内存（即使数据在L1/L2缓存中）。为此，我发现，L1高速缓存可以通过传递这些选项NVCC编译器跳过：如何将所有访问指向CUDA中的全局内存？

-Xptxas -dlcm=cg

CUDA文档状态这一点：

.cv Cache as volatile (consider cached system memory lines stale, fetch again).

所以，我假设当我要么运行 - dlcm = cg或-dlcm = cv，生成的PTX文件应该与正常生成的文件不同。

__global__ void rh_kernel(int *datainRowX, int *datainRowY) { 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    if (tid != 0) 
     return; 
    int i, x, y; 
    x = datainRowX[1]; 
    y = datainRowY[2]; 
    datainRowX[0] = x + y; 
} 

int main(int argc, char** argv) { 
    int* d_datainRowX; 
    cudaMalloc((void**)&d_datainRowX, sizeof(int) * 268435456); 

    int* d_datainRowY; 
    cudaMalloc((void**)&d_datainRowY, sizeof(int) * 268435456); 

    rh_kernel<<<1024, 1>>>(d_datainRowX, d_datainRowY); 
    cudaFree(d_datainRowX); cudaFree(d_datainRowY); 
    return(0); 
}

我注意到，无论选择我传递给NVCC编译器（“-Xptxas -dlcm = CG”或：

我的示例程序（该负载应与任一.cg或.CV追加） “-Xptxas -dlcm = cv”或者什么都不），在所有这三种情况下PTX生成的都是相同的。我正在使用-ptx选项来生成PTX文件。我错过了什么？有没有其他方法可以实现我正在做的事情？预先感谢您的时间。

来源

2017-01-16 Andrew Mathews

你在使用什么GPU（哪个cc）？ – pSoLT

我正在使用特斯拉K40 –

和nvcc版本6.5.12 –

根据Cuda Toolkit Documentation：在开普勒的GPU

L1高速缓存只保留本地存储器访问，诸如寄存器溢出和堆栈数据。全局负载仅缓存在L2 中（或在只读数据缓存中）。

基于GK110B产品，如特斯拉K40 GPU加速器，GK20A，和GK210默认保留这种行为

L1高速缓存未在全局内存默认使用开普勒读取。因此 - 当您添加-Xptxas -dlcm=cg时，PTX没有区别。

禁用L2缓存不是可能的。

来源

2017-01-16 09:12:09 pSoLT

谢谢你指出。为什么当我使用 -Xptxas -dlcm = cv 这不是用来直接访问全局内存吗？ –

.cv缓存为volatile（考虑缓存的系统内存行陈旧，再次获取）。 –

编辑：只要所有的访问都进入全局内存，我可以不禁用L2。因此，我想知道是否可以使用.cv –

如何将所有访问指向CUDA中的全局内存？

回答

相关问题