Performance-Problem mit dem Code CUDA, eine Variable nimmt 100-mal länger

Question

ich dies habe auszuführen: (es ist ein Chaos, weil ich jedes Bit bin zu ändern, um zu sehen, was falsch ist, und das, was der Compiler tut)

__device__ inline int f(int i, int j, int value) 
{ 
    int x; 
    int y; 
    int delta; 

    int* p = p_new_solution; 
    int pitch = p_new_solution_pitch; 

    int* p_row_i = (int*)((char*)p + i * pitch); 
    int p_i = p_row_i[threadIdx.x + blockIdx.x * blockDim.x]; 

    int* p_row_j = (int*)((char*)p + j * pitch); 
    int p_j = p_row_j[threadIdx.x + blockIdx.x * blockDim.x]; 

    delta = (tex2D(A_matrix, i, i) - tex2D(A_matrix, j, j)) * (tex2D(B_matrix, p_j, p_j) - tex2D(B_matrix, p_i, p_i)); 
    delta += (tex2D(A_matrix, i, j) - tex2D(A_matrix, j, i)) * (tex2D(B_matrix, p_j, p_i) - tex2D(B_matrix, p_i, p_j)); 


    for(int k = 0 ; k < n ; k++) 
    { 
     int* p_row = (int*)((char*)p + k * pitch); 
     int p_k = p_row[threadIdx.x + blockIdx.x * blockDim.x]; 

     int A_ki = tex2D(A_matrix, k, i); 
     int A_kj = tex2D(A_matrix, k, j); 
     int A_ik = tex2D(A_matrix, i, k); 
     int A_jk = tex2D(A_matrix, j, k); 
     int B_pkpj = tex2D(B_matrix, p_k, p_j); 
     int B_pkpi = tex2D(B_matrix, p_k, p_i); 
     int B_pjpk = tex2D(B_matrix, p_j, p_k); 
     int B_pipk = tex2D(B_matrix, p_i, p_k); 

     x = (A_ki - A_kj); 
     x *= (B_pkpj - B_pkpi); 

     y = (A_ik - A_jk); 
     y *= (B_pjpk - B_pipk); 

     x += y; 
    } 

    x -= ((tex2D(A_matrix, i, i) - tex2D(A_matrix, i, j)) * (tex2D(B_matrix, p_i, p_j) - tex2D(B_matrix, p_i, p_i))) + 
      ((tex2D(A_matrix, i, i) - tex2D(A_matrix, j, i)) * (tex2D(B_matrix, p_j, p_i) - tex2D(B_matrix, p_j, p_i))); 

    x -= ((tex2D(A_matrix, j, i) - tex2D(A_matrix, j, j)) * (tex2D(B_matrix, p_j, p_j) - tex2D(B_matrix, p_j, p_i))) + 
      ((tex2D(A_matrix, i, j) - tex2D(A_matrix, j, j)) * (tex2D(B_matrix, p_j, p_j) - tex2D(B_matrix, p_j, p_j))); 


    x += delta; 
    x *= 2; 

    return value; 
    //return x; 
} 

Das Problem ist mit diesen beiden Return-Anweisungen .. wenn ich return value, nimmt der ganze Kernel wie 300ms, wenn ich return x dauert es etwa 33000 ms. Was ist das Problem damit? Ich habe versucht, einige __syncthreads(), aber immer noch die gleiche schlechte Zeit.

Diese Return-Funktionen sind nicht der endgültige Code, ich brauche eine if else-Anweisung, um einen Rückgabewert zu wählen, entweder value oder value + x, und dies, wenn die else-Anweisung zu lange dauert.

Danke für jetzt.

Answer 1

Die Zeit, die Sie messen, ist nicht die Zeit, um eine Variable zurückzugeben, es ist die Zeit x zu berechnen. NVCC erkennt, dass Sie viel Code haben, der absolut nichts tut, da seine Ergebnisse nie verwendet werden, wenn Sie x nicht zurückgeben. Es entfernt den nutzlosen Code und macht die Funktion schneller.

Answer 2

return value gibt nur eines der Argumente der Funktion zurück und macht die gesamte Funktion zu einem No-Op. Ich schätze mal, es ist komplett weg optimiert. Wenn Sie return x es tut die eigentliche Arbeit und dauert 33s.