有沒有什麼辦法讓內核使用恆定或全局內存取決於數據大小

Question

__constant__ float constbuf[MAXSIZE] 
__device__ float *d_buf; 

__global__ void 
simple (float *buf2){ 
    //access buf2; 
} 


main(){ 

    int size, asize; 
    float *abuf, *d_buf2, *h_buf; 
    //... 
    if(size > MAXSIZE){ 
     cudaMalloc(&d_buf2, asize); 
     cudaMemcpy(d_buf2, h_buf, asize); 
     cudaMemcpyToSymbol(d_buf, &d_buf2, sizeof(d_buf2)); 
     cudaGetSymbolAddress((void **) &abuf, d_buf); 
    }else{ 
     cudaMemcpyToSymbol(constbuf, h_buf, asize); 
     cudaGetSymbolAddress((void **) &abuf, constbuf); 
    } 

    simple<<<grid, block, 0 ,stream>>>(abuf); 


} 

我想要做類似上面的事情，但我發現這樣內核沒有得到正確的緩衝區。無論如何要實現這一目標？我不想在內核如果可能的話

Answer 1

這個最好的解決辦法是有一個__device__內核完成大部分是包裹__device__內核的工作和兩個__global__內核加上「如果」條件。

例如：

__constant__ c_buf[MAXSIZE]; 

__device__ simple_core(float *buf, int len) 
{ 
// do something here. 
} 


__global__ simple_global_mem(float *d_buf, int len) 
{ 
    simple_core(d_buf, len); 
} 


__global__ simple_const_mem(int len) 
{ 
    simple_core(c_buf, len); 
} 

int main() 
{ 
// other code 

if (len < MAXSIZE) { 
    // cuda memcpy to symbol code here 
    simple_const_mem<<<threads, blocks>>>(len); 
} 
else { 
    simple_global_mem<<<threads, blocks>>>(d_buf, len): 
} 
}