阅读此问题后：“如何区分指向共享内存和全局内存的指针？” ，我决定试一试isspacep.local，isspacep.global并isspacep.shared在一个简单的测试程序中。
本地和共享内存的测试一直有效，但全局内存测试并不总是有效，例如在调试模式下编译设备代码时（-G）。
起初我以为编译器检测到我对全局内存使用了一个虚拟向量并对其进行了不同的处理，所以我使用了-Xcicc -O0 -Xptxas -O0（参见“完全禁用 NVCC 上的优化”）。如果我用 计算sm_30，则正确检测到全局内存。但是，如果我使用sm_20or进行计算，sm_21则不会检测到全局内存。请注意，使用-G, 任何sm >= 20作品。
我在这里缺少什么吗？使用时是否有额外的标志给予编译器-G可以解释这些差异？
汇编
nvcc test_pointer.cu -arch=sm_20 -Xcicc -O0 -Xptxas -O0 -Xptxas -v -o test_pointer
代码
#include <stdio.h>
#include <cuda.h>
#define CUDA_CHECK_ERROR() __cuda_check_errors(__FILE__, __LINE__)
#define CUDA_SAFE_CALL(err) __cuda_safe_call(err, __FILE__, __LINE__)
inline void __cuda_check_errors(const char *filename, const int line_number)
{
cudaError err = cudaDeviceSynchronize();
if(err != cudaSuccess)
{
printf("CUDA error %i at %s:%i: %s\n",
err, filename, line_number, cudaGetErrorString(err));
exit(-1);
}
}
inline void __cuda_safe_call(cudaError err, const char *filename, const int line_number)
{
if (err != cudaSuccess)
{
printf("CUDA error %i at %s:%i: %s\n",
err, filename, line_number, cudaGetErrorString(err));
exit(-1);
}
}
__device__ unsigned int __isLocal(const void *ptr)
{
unsigned int ret;
asm volatile ("{ \n\t"
" .reg .pred p; \n\t"
" isspacep.local p, %1; \n\t"
" selp.u32 %0, 1, 0, p; \n\t"
#if (defined(_MSC_VER) && defined(_WIN64)) || defined(__LP64__)
"} \n\t" : "=r"(ret) : "l"(ptr));
#else
"} \n\t" : "=r"(ret) : "r"(ptr));
#endif
return ret;
}
__device__ unsigned int __isShared(const void *ptr)
{
unsigned int ret;
asm volatile ("{ \n\t"
" .reg .pred p; \n\t"
" isspacep.shared p, %1; \n\t"
" selp.u32 %0, 1, 0, p; \n\t"
#if (defined(_MSC_VER) && defined(_WIN64)) || defined(__LP64__)
"} \n\t" : "=r"(ret) : "l"(ptr));
#else
"} \n\t" : "=r"(ret) : "r"(ptr));
#endif
return ret;
}
__device__ void analyze_pointer(const void *ptr)
{
printf("\t* is local: %u\n", __isLocal(ptr));
printf("\t* is global: %u\n", __isGlobal(ptr));
printf("\t* is shared: %u\n", __isShared(ptr));
}
template <typename T, unsigned int N>
__global__ void test_kernel(T *vec)
{
// Shared array
__shared__ T shared_vec[10];
// Register array
T reg[10];
if (blockIdx.x == 0 && threadIdx.x == 0)
{
printf("Register array:\n");
analyze_pointer(®);
printf("\nGlobal array:\n");
analyze_pointer(vec);
printf("\nShared array:\n");
analyze_pointer(&shared_vec);
}
}
int main()
{
typedef float type_t;
const unsigned int N = 128;
type_t* d_vec;
CUDA_SAFE_CALL(cudaMalloc(&d_vec, N * sizeof(type_t)));
test_kernel<type_t, N><<<1, N>>>(d_vec);
CUDA_CHECK_ERROR();
CUDA_SAFE_CALL(cudaFree(d_vec));
}
输出
Register array:
* is local: 1
* is global: 0
* is shared: 0
Global array:
* is local: 0
* is global: 0 (or 1 with -G or sm_30)
* is shared: 0
Shared array:
* is local: 0
* is global: 0
* is shared: 1
硬件/软件属性
在 Arch Linux 64 位上使用 CUDA 5.0、GeForce GT 650M (CC 3.0)、驱动程序 319.17 进行了测试。
更新#1
我刚刚使用带有 304.88 驱动程序的 Tesla C2070 (CC 2.0)、Linux 64 位上的 CUDA 5.0 测试了此代码，并且它可以工作。当优化关闭时，即检测到全局内存-arch=sm_20 -Xcicc -O0，或者添加额外printf("\t* ptr = %ld\n", ptr);的（参见@RobertCrovella 的评论）。听起来确实像驱动程序问题。
更新#2
我做了一些更多的测试，下面是我使用 CC 3.0 设备得到的结果，具体取决于我编译程序的方式：
-arch=sm_30 ---> undetected (probably optimized)
-arch=sm_30 -Xcicc -O0 -Xptxas -O0 ---> OK
-arch=sm_30 -G ---> OK
-arch=compute_30 -code=sm_30 -Xcicc -O0 -Xptxas -O0 ---> OK
-arch=compute_30 -code=sm_30 -G ---> OK
-arch=compute_30 -code=compute_30 -Xcicc -O0 -Xptxas -O0 ---> undetected
-arch=compute_30 -code=compute_30 -G ---> OK
-arch=sm_20 ---> undetected
-arch=sm_20 -Xcicc -O0 -Xptxas -O0 ---> undetected
-arch=sm_20 -G ---> OK
-arch=compute_20 -Xcicc -O0 -Xptxas -O0 ---> undetected
-arch=compute_20 -G ---> OK
-arch=compute_20 -code=sm_20 -Xcicc -O0 -Xptxas -O0 ---> runtime error (as expected)
-arch=compute_20 -code=sm_20 -G ---> runtime error (as expected)
-arch=compute_20 -code=compute_20 -Xcicc -O0 -Xptxas -O0 ---> undetected
-arch=compute_20 -code=compute_20 -G ---> OK
-arch=compute_20 -code=sm_30 ---> undetected (probably optimized)
-arch=compute_20 -code=sm_30 -Xcicc -O0 -Xptxas -O0 ---> OK
-arch=compute_20 -code=sm_30 -G ---> OK