我可以檢查地址是否在共享內存中？

我想寫以下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通常可以知道它們是用於共享內存還是不。

來源

2017-02-28 einpoklum

您可以通過內嵌位的「集結號」使用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 
    } 
}

來源

2017-03-01 08:21:54 tera

偉大的:-)你能解釋一下爲什麼我一開始需要額外的Win64專用伏都教嗎？另外，PTX指南說我們有「通用」存儲器地址，據說是空間專用地址。這對你的答案如何？ – einpoklum

我補充說，「巫術」，以防萬一任何人試圖在32位模式下編譯。如果你知道你的代碼總是64位的，你可以直接使用指針的''l''約束。 – tera

...編輯到你的答案。 – einpoklum

_{這是@萬億的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位代碼，更換"l"約束（一64位地址）與"r"。

來源

2017-03-01 09:38:22 einpoklum

你試過這個嗎？ – tera

@tera：Fankly - 只編譯它，我沒有在這臺GPU上使用第二臺計算機...... – einpoklum

好吧，我試過了，但它沒有通過ptxas。使用'isspace'，我得到'不是任何已知指令的名稱：'isspace''，'isspacep'如預期的''參數不匹配指令'isspacep'。 – tera

我可以檢查地址是否在共享內存中？

回答

相關問題