CUDA全局（如C語言）分配給設備內存的動態數組

Question

所以，我試圖編寫一些利用Nvidia CUDA架構的代碼。我注意到，複製到設備和從設備複製真的傷害了我的整體性能，所以現在我試圖將大量數據移動到設備上。由於這些數據被用於多種功能，我希望它是全球性的。是的，我可以通過指針，但我真的很想知道在這種情況下如何使用全局變量。

所以，我有設備功能，要訪問設備分配數組。

理想情況下，我可以這樣做：

__device__ float* global_data; 

main() 
{ 
    cudaMalloc(global_data); 
    kernel1<<<blah>>>(blah); //access global data 
    kernel2<<<blah>>>(blah); //access global data again 
} 

不過，我還沒有想出如何創建一個動態數組。我想出了一個變通通過聲明數組如下：

__device__ float global_data[REALLY_LARGE_NUMBER]; 

儘管這並不需要cudaMalloc電話，我寧願動態分配方法。

Answer 1

像這樣的東西應該可能工作。

#include <algorithm> 

#define NDEBUG 
#define CUT_CHECK_ERROR(errorMessage) do {         \ 
     cudaThreadSynchronize();           \ 
     cudaError_t err = cudaGetLastError();        \ 
     if(cudaSuccess != err) {           \ 
        fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \ 
              errorMessage, __FILE__, __LINE__, cudaGetErrorString(err));\ 
        exit(EXIT_FAILURE);             \ 
       } } while (0) 


__device__ float *devPtr; 

__global__ 
void kernel1(float *some_neat_data) 
{ 
    devPtr = some_neat_data; 
} 

__global__ 
void kernel2(void) 
{ 
    devPtr[threadIdx.x] *= .3f; 
} 


int main(int argc, char *argv[]) 
{ 
    float* otherDevPtr; 
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr)); 
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr)); 

    kernel1<<<1,128>>>(otherDevPtr); 
    CUT_CHECK_ERROR("kernel1"); 

    kernel2<<<1,128>>>(); 

    CUT_CHECK_ERROR("kernel2"); 

    return 0; 
} 

給它一個旋風。

Answer 2

花一些時間專注於NVIDIA提供的豐富文檔。

從編程指南：

float* devPtr; 
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr)); 
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr)); 

這是一個如何分配內存簡單的例子。現在，在你的內核，你應該接受一個指向浮動，像這樣：

__global__ 
void kernel1(float *some_neat_data) 
{ 
    some_neat_data[threadIdx.x]++; 
} 

__global__ 
void kernel2(float *potentially_that_same_neat_data) 
{ 
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f; 
} 

所以現在你可以調用它們像這樣：

float* devPtr; 
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr)); 
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr)); 

kernel1<<<1,128>>>(devPtr); 
kernel2<<<1,128>>>(devPtr); 

由於這些數據在衆多 使用功能，我想它是全球性的 。

使用全局變量的原因很少。這絕對不是一個。我將把它作爲一個練習來擴展這個例子，包括將「devPtr」移動到全局範圍。

編輯：

好，最根本的問題是這樣的：你的內核只能接入設備內存以及他們能夠使用的唯一的全球範圍的指針是GPU的。當從CPU調用內核時，幕後會發生什麼，指針和原語在內核執行之前被複制到GPU寄存器和/或共享內存中。

所以我可以建議的最接近的是：使用cudaMemcpyToSymbol（）來實現你的目標。但是，在背景中，考慮一種不同的方法可能是正確的。

#include <algorithm> 

__constant__ float devPtr[1024]; 

__global__ 
void kernel1(float *some_neat_data) 
{ 
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1]; 
} 

__global__ 
void kernel2(float *potentially_that_same_neat_data) 
{ 
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2]; 
} 


int main(int argc, char *argv[]) 
{ 
    float some_data[256]; 
    for (int i = 0; i < sizeof(some_data)/sizeof(some_data[0]); i++) 
    { 
     some_data[i] = i * 2; 
    } 
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr))); 
    float* otherDevPtr; 
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr)); 
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr)); 

    kernel1<<<1,128>>>(otherDevPtr); 
    kernel2<<<1,128>>>(otherDevPtr); 

    return 0; 
} 

對於這個例子，不要忘記'--host-compilation = C++'。

Answer 3

呃，這正是將devPtr移動到全局範圍的問題，這是我的問題。

我有一個實現，確實如此，兩個內核有一個指向數據傳入的指針。我明確不想傳入這些指針。

我已經仔細閱讀了文檔，打了nvidia論壇（和谷歌搜索了一個小時左右），但我還沒有找到一個實際運行的全球動態設備數組的實現（我試過幾個編譯，然後以新的和有趣的方式失敗）。

Answer 4

查看SDK附帶的示例。這些示例項目中的很多都是一個體面的學習方式。

Answer 5

我繼續嘗試分配臨時指針並將它傳遞給類似於kernel1的簡單全局函數的解決方案。

的好消息是，它的工作:)

不過，我認爲它混淆編譯器，我現在得到「諮詢：不能說什麼指針指向，假設全球內存空間」每當我試圖訪問全球數據。幸運的是，這個假設是正確的，但警告很煩人。

無論如何，爲了記錄 - 我已經看過很多例子，並且通過nvidia練習運行，其中的重點是讓輸出說「正確！」。不過，我還沒有看過他們的全部。如果有人知道他們做動態全局設備內存分配的sdk示例，我仍然想知道。

Answer 6

由於此數據用於衆多功能，我希望它是全球性的。

-

有兩種使用全局幾個很好的理由。這絕對不是一個。我將把它作爲 練習來擴展此示例，以便將「devPtr」移動到全局範圍。

如果內核在由陣列的大常量結構工作是什麼？使用所謂的常量內存不是一種選擇，因爲它的大小非常有限......所以你必須把它放在全局內存中。