アドレスが共有メモリにあるかどうかを確認できますか？

Question

私は、次のCUDA機能書きたい：ホスト側で

void foo(int* a, size_t n) 
{ 
    if (/* MAGIC 1 */) { 
     // a is known to be in shared memory, 
     // so use it directly 
    } 
    else { 
     // make a copy of a in shared memory 
     // and use the copy 
    } 
} 

を、私たちは、ポインタは、デバイスのメモリにあるかかどうかを私たちに伝えることができた、cudaPointerGetAttributesの形でやや関連の施設を持っていますホストメモリ。おそらく、デバイスコード内のポインタを区別する方法もありますし、グローバルポインタからの共有を識別することもできます。あるいは、おそらくそれを行うためのコンパイル時のメカニズムがあるかもしれません。なぜなら、デバイスの機能はカーネルにコンパイルされ、自立していないからです。したがって、nvccは共有メモリで使用されているかどうか、ない。

Answer 1

あなたはインライン "組立" のビットを経由してisspacep PTX instructionを使用することができます。

// First, a pointer-size-related definition, in case 
// this code is being compiled in 32-bit rather than 
// 64-bit mode; if you know the code is always 64-bit 
// you can just use the "l" 

#if defined(_WIN64) || defined(__LP64__) 
# define PTR_CONSTRAINT "l" 
#else 
# define PTR_CONSTRAINT "r" 
#endif 

__device__ int isShared(void *ptr) 
{ 
    int res; 
    asm("{" 
     ".reg .pred p;\n\t" 
     "isspacep.shared p, %1;\n\t" 
     "selp.b32 %0, 1, 0, p;\n\t" 
     "}" : 
     "=r"(res): PTR_CONSTRAINT(ptr)); 
    return res; 
} 

をので、あなたの例では、

void foo(int* a, size_t n) 
{ 
    if (isShared(a)) { 
     // a is known to be in shared memory, 
     // so use it directly 
    } else { 
     // make a copy of a in shared memory 
     // and use the copy 
    } 
} 

Answer 2

_{なりこれは、@ TERAのanswerを一般化したものです。すべての可能なオンデバイスのメモリ空間のための同様の機能を定義して、次のコードから}

使用is_in_shared_memory()は：

#ifndef STRINGIFY 
#define STRINGIFY(_q) #_q 
#endif 

#define IS_IN_MEMORY_SPACE(_which_space) \ 
__forceinline__ __device__ int is_in_ ## _which_space ## _memory (const void *ptr) \ 
{ \ 
    int result; \ 
    asm ("{" \ 
     ".reg .pred p;\n\t" \ 
     "isspacep." STRINGIFY(_which_space) " p, %1;\n\t" \ 
     "selp.b32 %0, 1, 0, p;\n\t" \ 
     "}" \ 
     : "=r"(result) : "l"(ptr)); \ 
    return result; \ 
} 

IS_IN_MEMORY_SPACE(const) 
IS_IN_MEMORY_SPACE(global) 
IS_IN_MEMORY_SPACE(local) 
IS_IN_MEMORY_SPACE(shared) 

#undef IS_IN_MEMORY_SPACE 

あなたは32ビットコードを構築している場合は、64（"l"制約を交換しますビットアドレス）を"r"に設定します。