2013-03-07 114 views
0

我用CUDA開發了以下插值,我正在尋找一種改進插值的方法。出於某些原因,我不想使用CUDA紋理。優化CUDA插值

另一點,我已經注意到,由於某些未知的原因,插值不是在整個向量上執行,在我的情況下,如果向量的大小優於線程的數量(例如向量。大小1000,等於512多個線程的,.線程做它的第一份工作,這一切我想優化singleInterp功能

這裏是我的代碼:

__device__ float singleInterp(float* data, float x, int lx_data) { 

float res = 0;  
int i1=0; 
int j=lx_data; 
int imid; 

while (j>i1+1) 
{ 
    imid = (int)(i1+j+1)/2; 
    if (data[imid]<x) 
     i1=imid; 
    else 
     j=imid; 
} 
if (i1==j) 
    res = data[i1+lx_data]; 
else 
    res =__fmaf_rn(__fdividef(data[j+lx_data]-data[i1+lx_data],(data[j]-data[i1])),x-data[i1], data[i1+lx_data]); 

return res; 

} 

內核:

__global__ void linearInterpolation(float* data, float* x_in, int lx_data) { 

int i = threadIdx.x + blockDim.x * blockIdx.x; 
int index = i; 
if (index < lx_data) 
    x_in[index] = singleInterp(data, x_in[index], lx_data); 
} 
+1

您啓動了多少個區塊?你能列出你的內核設置和調用代碼嗎?您需要正確計算需要處理的塊的數量,以及每個塊有多少個線程。在你上面的例子中,你需要2塊。 – lmortenson 2013-03-07 14:48:06

回答

1

看來你對1D線性插值感興趣。我已經有了優化這樣一種插值的問題,我結束了與下面的代碼

__global__ void linear_interpolation_kernel_function_GPU(double* __restrict__ result_d, const double* __restrict__ data_d, const double* __restrict__ x_out_d, const int M, const int N) 
{ 
    int j = threadIdx.x + blockDim.x * blockIdx.x; 

    if(j<N) 
    { 
     double reg_x_out = x_out_d[j/2]+M/2; 
     int k = floor(reg_x_out); 
     double a = (reg_x_out)-floor(reg_x_out); 
     double dk = data_d[2*k+(j&1)]; 
     double dkp1 = data_d[2*k+2+(j&1)]; 
     result_d[j] = a * dkp1 + (-dk * a + dk); 
    } 
} 

假定該數據在整數節點-M/2M/2之間進行採樣。 該代碼與1D紋理插值「等效」,如以下web-page所述。對於1D線性紋理插值,請參見CUDA-Programming-Guide的圖13。有關不同解決方案之間的比較,請參閱以下thread