一段时间以来,我一直在为一个似乎无法找到解决方案的问题而苦苦挣扎。
问题是,当我尝试在 Visual Studio 2008 下使用 Nvidia Nsight 调试我的 CUDA 代码时,我在使用共享内存时得到了奇怪的结果。
我的密码是:
template<typename T>
__device__
T integrate()
{
extern __shared__ T s_test[]; // Dynamically allocated shared memory
/**** Breakpoint (1) here ****/
int index = threadIdx.x + threadIdx.y * blockDim.x; // Local index in block. Column major ordering
if(index < 64 && blockIdx.x==0) { // Only work on a few values. Just testing
s_test[index] = (T)index;
/* Some other irelevant code here */
}
return v;
}
当我到达
Breakpoint (1)
并检查 Visual Studio Watch 窗口中的共享内存时,只有数组的前 8 个值发生变化,其他值保持为空。我希望所有前 64 个都这样做。
我认为这可能与所有扭曲未同时执行有关。所以我尝试同步它们。我在里面添加了这段代码
integrate()
template<typename T>
__device__
T integrate()
{
/* Old code is still here */
__syncthreads();
/**** Breakpoint (2) here ****/
if(index < 64 && blockIdx.x==0) {
T tmp = s_test[index]; // Write to tmp variable so I can inspect it inside Nsight Watch window
v = tmp + index; // Use `tmp` and `index` somehow so that the compiler doesn't optimize it out of existence
}
return v;
}
但问题依旧。此外,
tmp
内的其余值不是 VS 的 Watch 窗口所指示的 0
。
我必须提到,跨过
__syncthreads()
需要很多步,所以当我到达它时,我只是跳到Breakpoint (2)
。这是怎么回事?
编辑有关系统/启动配置的信息
系统
设备 GeForce 9500 GT
IDE
编译器命令
1> "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\\bin\nvcc.exe" -G -gencode=arch=compute_10,code=\"sm_10,compute_10\" --machine 32 -ccbin "C:\Program Files\Microsoft Visual Studio 9.0\VC\bin" -D_NEXUS_DEBUG -g -D_DEBUG -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd " -I"inc" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\\include" -maxrregcount=0 --compile -o "Debug/process_f2f.cu.obj" process_f2f.cu
启动配置。共享内存大小似乎并不重要。我试过几个版本。我合作最多的是:
您是否尝试过在分配值后放置
__syncthreads()
?
template<typename T>
__device__
T integrate()
{
extern __shared__ T s_test[]; // Dynamically allocated shared memory
int index = threadIdx.x + threadIdx.y * blockDim.x; // Local index in block. Column major ordering
if(index < 64 && blockIdx.x==0) { // Only work on a few values. Just testing
s_test[index] = (T)index;
/* Some other irelevant code here */
}
__syncthreads();
/**** Breakpoint (1) here ****/
return v;
}
并尝试查看此断点处的值。