この質問を読んだ後: 「共有メモリとグローバルメモリへのポインタを区別する方法は?」、私は試してみることにしましたisspacep.local、isspacep.globalそしてisspacep.shared簡単なテストプログラムで。
ローカル メモリと共有メモリのテストは常に機能しますが、グローバル メモリ テストは常に機能するとは限りません。たとえば、デバイス コードがデバッグ モードでコンパイルされた場合に機能します ( -G)。
最初は、コンパイラがグローバルメモリにダミーベクトルを使用していることを検出し、別の方法で処理したと思ったので、使用しました-Xcicc -O0 -Xptxas -O0(cf. 「NVCC の最適化を完全に無効にする」 )。で計算するとsm_30、グローバル メモリが正しく検出されます。sm_20ただし、またはで計算すると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