在不使用內核的情況下在CUDA中寫入共享內存

我想在main（）函數中創建一個數組，輸入所有正確的值，然後讓該數組立即可以被共享內存中的線程使用。在不使用內核的情況下在CUDA中寫入共享內存

我在CUDA中查找每個使用共享內存的例子都有線程在寫入共享數組，但我希望我的共享數組在內核啓動之前立即可用。

任何幫助做到這一點將不勝感激。提前致謝！

某些上下文：我想要的共享數組永遠不會更改，並且會被所有線程讀取。

編輯：顯然這是不可能與共享內存。有誰知道是否可以使用只讀緩存？

來源

2014-06-25 Jordan

這是不可能的。填充shared memory的唯一方法是使用CUDA內核中的線程。

如果您希望在啓動時可以使用一組（只讀）數據供內核使用，則當然可以使用__constant__ memory。這樣的存儲器可以使用文檔中指示的API在主機代碼上設置，即cudaMemcpyToSymbol。

__constant__只有當每個線程在給定訪問週期（例如，訪問週期）訪問相同位置時，存儲器才真正有用。

int myval = constant_data[12];

否則使用普通全局內存，無論是靜態還是動態分配，使用適當的主機API來初始化（動態：cudaMemcpy，靜態：cudaMemcpyToSymbol）。

來源

2014-06-25 20:11:30

太臭了。你知道是否有可能用compute 3.5設備中的只讀緩存來實現我想要的嗎？謝謝。 – Jordan

讓我們清楚，SO不是聊天服務。如果您有新問題，請將其作爲新問題發佈。 –

如果60％的線程正在讀取相同的地址，而其他40％的線程讀取到3個其他地址的比例相同，您是否認爲常量內存比全局內存更有效？ – Jordan

雖然你所要求的特定行爲自動是不可能的，這其實是一個相當普遍的CUDA範式：

首先，擁有所有線程複製表到SHMEM。

同步線程

訪問內核中的數據。

如果您對數據進行了相當隨機的訪問，並且您希望平均觸及每個條目超過幾次，這可能會帶來巨大的性能提升。基本上你使用shmem作爲一個託管緩存，並將來自DRAM的負載聚合成shmem一次，以便多次使用。此外，shmem對於未聚合的負載沒有懲罰。

例如，你可以編寫它是這樣的：

const int buffer_size = 8192; // assume an 8k buffer 
float *device_buffer = ; // assume you have a buffer already on the device with the data you want. 

my_kernel<<<num_blocks, num_threads, buffer_size>>>(..., buffer_size, device_buffer); 

__global__ void my_kernel(..., int buffer_size, const float *device_buffer) { 
    extern __shared__ float shmem_buffer[]; 
    for (int idx = threadIdx.x; idx < buffer_sze; idx += blockDim.x) { 
     shmem_buffer[idx] = device_buffer[idx]; 
    } 
    __syncthreads(); 

    // rest of your kernel goes here. You can access data in shmem_buffer; 
}

換句話說，你只需要明確的代碼副本。由於來自DRAM的所有負載都將完美結合，因此這應該接近最佳效率。

來源

2014-06-26 20:01:33

在不使用內核的情況下在CUDA中寫入共享內存

回答

相關問題