PyCuda를 이용한 유전자 세포 자동화, CUDA 커널에 셀당 많은 데이터를 효율적으로 전달하는 방법은 무엇입니까?

저는 PyCuda를 사용하여 유전자 세포 자동화를 개발 중입니다. 각 셀에는 셀 매개 변수와 함께 많은 게놈 데이터가 있습니다. 1) 셀 데이터를 CUDA 커널에 전달한 다음 2)이 데이터를 처리하는 가장 효율적인 방법이 무엇인지 궁금합니다.PyCuda를 이용한 유전자 세포 자동화, CUDA 커널에 셀당 많은 데이터를 효율적으로 전달하는 방법은 무엇입니까?

필자는 특히 나쁜 해결책으로 시작했지만 여전히 해결책이되고 있습니다. 그것은 각 매개 변수를 별도의 배열로 전달한 다음 switch-case와 많은 중복 코드로 처리합니다.

그런 다음 커널 기능 당 매우 많은 수의 매개 변수로 신속하게 끝낼 수 있고 다시 작성하기로 결정했습니다.

둘째 해결 방법은 추가 매개 변수가있는 단일 배열에 셀의 모든 매개 변수를 저장하는 것이 었습니다. 훨씬 더 우아한 코드 였지만 놀랍게도 코드는 10 배 느리게 실행됩니다! 셀의 현재의 '맛', 대량 - 배 (INT) :

(FC, 맥, TC) :

이 좀 더 명확하게하기 위해, 내가 필요로하는 데이터의 전체 목록은 셀당 저장하기 온도
(RFC, RMC, RTC) 3X (INT) - 셀의 현재 레지스터
각각 인접 대 (FI, MI, 티타늄) * 8 배 (INT) - 입력 값
(RFI , Rmi, Rti) : 8 * 3x (int) - 들어오는 값
게이트 방향 : 1x (UCHAR)
실행 포인터 : 1 배 (UCHAR)
현재 마이크로 작업 메모리 : 32 배속 (UCHAR)
마지막 단계의 마이크로 작업 메모리 : 32 배속 (UCHAR)

내가 분할 해요 2 단계의 자동화 단계. 첫 번째 (방출 단계)는 각 셀의 이웃에 대해 (Fi, Mi, Ti)를 계산합니다. 두 번째 (흡수 단계)는 현재 셀의 상태와 8x (Fi, Mi, Ti) 값을 블렌딩합니다. 게놈이나 레지스터가 아직 구현되지 않았지만 데이터를 전달해야합니다.

Mk = 64 
Tk = 1000 
CELL_LEN = 120 # number of parameters per cell 

emit_gpu = ElementwiseKernel("int *cells, int w, int h", """ 
    int x = i/h; 
    int y = i % h; 
    int ii = i * CN; 

    int Fc = cells[ii]; 
    int Mc = cells[ii+1]; 
    int Tc = cells[ii+2]; 
    float M = (float) Mc; 
    float T = (float) Tc; 
    int Mi = (int) (fmin(1, T/Tk) * M); 
    cells[ii+1] = Mc - Mi; 
    cells[ii+2] = Tc - (int) (T * fmin(1, T/Tk)); 

    int Mbase = Mi/8; 
    int Mpart = Mi % 8; 
    int Madd; 
    int iii, xo, yo; 

    for (int cc = 0; cc < 9; cc++) { 
     int c = (cc + Fc) % 9; 
     if (c == 4) continue; 
     xo = x + c%3 - 1; 
     if (xo < 0) xo = w + xo; else if (xo >= w) xo = xo - w; 
     yo = y + c/3 - 1; 
     if (yo < 0) yo = h + yo; else if (xo >= w) yo = yo - h; 
     if (Mpart > 0) { Madd = 1; Mpart--;} else Madd = 0; 
     if (c > 4) c--; 
     iii = (xo * h + yo) * CN + 6 + c*3; 

     cells[iii] = Fc; 
     cells[iii+1] = Mbase + Madd; 
     cells[iii+2] = Tc; 

    } 
""", "ca_emit", preamble=""" 
#define Tk %s 
#define CN %s 
""" % (Tk, CELL_LEN)) 

absorb_gpu = ElementwiseKernel("int *cells, int *img, int w, int h", """ 
    int ii = i * CN; 
    int Fc = cells[ii]; 
    int Mc = cells[ii+1]; 
    int Tc = cells[ii+2]; 

    for (int c=0; c < 8; c++){ 
     int iii = ii + c * 3 + 6; 
     int Fi = cells[iii]; 
     int Mi = cells[iii+1]; 
     int Ti = cells[iii+2]; 

     int dF = Fi - Fc; 
     if (dF > 180) Fc += 360; 
     if (dF < -180) Fc -= 360; 
     float sM = Mi + Mc; 
     if (sM != 0) sM = Mi/sM; else sM = 0; 
     dF = (int) (Fi - Fc) * sM; 
     int dM = Mi; 
     int dT = fabs((float) (Fi - Fc)) * fmin((float) Mc, (float) Mi)/Mk + (Ti - Tc) * sM; 
     Fc += dF; 
     Mc += dM; 
     Tc += dT; 
     Fc = Fc % 360; 
     if (Fc < 0) Fc += 360; 
     if (Tc > Tk) Tc = Tk; 
    }  

    cells[ii] = Fc; 
    cells[ii+1] = Mc; 
    cells[ii+2] = Tc; 
    cells[ii+18] = (cells[ii+18] + 1) % 8; 

    img[i] = hsv2rgb(Fc, Tc, Mc); 

""", "ca_absorb", preamble=""" 
#include <math.h> 
#define Mk %s 
#define Tk %s 
#define CN %s 

__device__ uint hsv2rgb(int hue, int sat, int val) { 
    // skipped for brevity 
} 
""" % (Mk, Tk, CELL_LEN))

두 변종은 정확히 같은 CA의 동작을 생산하지만, 후자는 훨씬 느린 실행 :

Mk = 64 
Tk = 1000 

emit_gpu = ElementwiseKernel("int3 *cells, int3 *dcells0, int3 *dcells1, int3 *dcells2, int3 *dcells3, int3 *dcells4, int3 *dcells5, int3 *dcells6, int3 *dcells7, int w, int h", """ 
    int x = i/h; 
    int y = i % h; 

    int3 cell = cells[i]; 
    float M = (float) cell.y; 
    float T = (float) cell.z; 
    int Mi = (int) (fmin(1, T/Tk) * M); 
    cells[i].y -= Mi; 
    cells[i].z -= (int) (T * fmin(1, T/Tk)/1); 

    int Fi = cell.x; 
    int Mbase = Mi/8; 
    int Mpart = Mi % 8; 
    int Madd; 
    int Ti = cell.z; 
    int ii, xo, yo; 

    for (int cc = 0; cc < 9; cc++) { 
     int c = (cc + Fi) % 9; 
     if (c == 4) continue; 
     xo = x + c%3 - 1; 
     if (xo < 0) xo = w + xo; 
     if (xo >= w) xo = xo - w; 
     yo = y + c/3 - 1; 
     if (yo < 0) yo = h + yo; 
     if (xo >= w) yo = yo - h; 
     ii = xo * h + yo; 
     if (Mpart > 0) { Madd = 1; Mpart--;} else Madd = 0; 
     switch(c) { 
     case 0: dcells0[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 1: dcells1[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 2: dcells2[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 3: dcells3[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 5: dcells4[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 6: dcells5[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 7: dcells6[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     case 8: dcells7[ii] = make_int3(Fi, Mbase + Madd, Ti); break; 
     default: break; 
     } 

    } 
""", "ca_prepare", preamble=""" 
#define Tk %s 
""" % Tk) 

absorb_gpu = ElementwiseKernel("int3 *cells, int3 *dcells0, int3 *dcells1, int3 *dcells2, int3 *dcells3, int3 *dcells4, int3 *dcells5, int3 *dcells6, int3 *dcells7, int *img, int w, int h", """ 
    int3 cell = cells[i]; 

    int3 dcell = dcells0[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 

    dcell = dcells1[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    dcell = dcells2[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    dcell = dcells3[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    dcell = dcells4[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    dcell = dcells5[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    dcell = dcells6[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    dcell = dcells7[i]; 
    cell = cell + calc_d(cell.x, cell.y, cell.z, dcell.x, dcell.y, dcell.z); 
    cell.x = cell.x % 360; 
    if (cell.x < 0) cell.x += 360; 
    if (cell.z > Tk) cell.z = Tk; 

    cells[i] = cell; 
    img[i] = hsv2rgb(cell); 

""", "ca_calc", preamble=""" 
#include <math.h> 
#define Mk %s 
#define Tk %s 

__device__ int3 operator+(const int3 &a, const int3 &b) { 
    return make_int3(a.x+b.x, a.y+b.y, a.z+b.z); 
} 

__device__ int3 calc_d(int Fc, int Mc, int Tc, int Fi, int Mi, int Ti) { 
    int dF = Fi - Fc; 
    if (dF > 180) Fc += 360; 
    if (dF < -180) Fc -= 360; 
    float sM = Mi + Mc; 
    if (sM != 0) sM = Mi/sM; else sM = 0; 
    dF = (int) (Fi - Fc) * sM; 
    int dM = Mi; 
    int dT = fabs((float) (Fi - Fc)) * fmin((float) Mc, (float) Mi)/Mk + (Ti - Tc) * sM; 
    return make_int3(dF, dM, dT); 
} 

__device__ uint hsv2rgb(int3 pixel) { 
    // skipped for brevity 
} 
""" % (Mk, Tk, RAM))

두 번째와 현재 솔루션 :

그래서, 내 최초의 솔루션의 코드였다.

GTX 타이탄 :

필드의 크기 : 1900x1080 세포
솔루션 # 1 : ~ 200 단계/s의
솔루션 # 2 : ~ 20 단계/s의

GT의 630M :

필드 크기 : 1600x900 셀
해결책 # 1 : ~ 7.8 단계/s의
솔루션 # 2 : ~ 1.5 단계/s의

당신이 필요로하는 경우에 '두 솔루션 놀 주시기 바랍니다 :

Solution #1 full source

Solution #2 full source

하나를 단서 또는 조언을 환영합니다 :

왜 성능이 느려 집니까?
솔루션 # 2의 성능을 적어도 # 1 수준까지 올릴 수 있습니까?
다른 해결책이 더 좋을까요?

출처

2014-12-16 a5kin

2688 CUDA 코어를 가진 카드가 96 코어를 가진 카드에 비해 놀랍습니다. – AlexanderBrevig

죄송합니다. 질문을 잘못 읽었습니다 - 해결책 1 : 2, 카드 1 : 2와 관련이 없습니다. 죄송합니다. – AlexanderBrevig

바로, 최적화를 시작한 첫 번째 문제는 왜 두 가지 솔루션간에 성능이 크게 과장되었는지를 아는 것입니다. 다른 카드 결과는 비교하기 쉽습니다. 코어가 적 으면 속도가 느려집니다. GT630에서 5 배 느린 타이탄 VS에서 10 배 느립니다. – a5kin

좋아요, 나는 두 번째 솔루션을 거의 15 배 빠르게 실행하는 방법을 관리했습니다. 다음과 같은 변경이 이루어졌다 :

는 int에서 int4에 주요 매개 변수의 배열을 변환합니다. 이로 인해 int3의 솔루션보다 훨씬 빨랐습니다. 그러나 여분의 공간은 사용되지 않은 상태로 남아 있습니다 (.w 치수). [3 배속 향상]

WIDTHxHEIGHT 그룹의 Repack 관련 매개 변수입니다. 따라서 모양이 (너비, 높이, N)에서 (N, 너비, 높이)로 변경되었습니다. 그룹 내의 요소들은 함께 처리되는 경향이 있으므로 메모리 액세스가보다 효율적으로 이루어졌습니다. 재 포장에 대한 단서를 공원 젊은-배에

Mk = 64 
Tk = 1000 

emit_gpu = ElementwiseKernel("int4 *cells, int w, int h, int cn", """ 
    int x = i/h; 
    int y = i % h; 

    int4 cell = cells[i]; 
    int Fc = cell.x; 
    int Mc = cell.y; 
    int Tc = cell.z; 
    float M = (float) Mc; 
    float T = (float) Tc; 
    int Mi = (int) (fmin(1, T/Tk) * M); 
    cells[i] = make_int4(Fc, Mc - Mi, Tc - (int) (T * fmin(1, T/Tk)), 0); 

    int Mbase = Mi/8; 
    int Mpart = Mi % 8; 
    int Madd; 
    int ii; 
    int xo, yo; 

    int cnn = 0; 
    for (int dx = -1; dx < 2; dx++) { 
     xo = x + dx; 
     if (xo < 0) xo = w + xo; else if (xo >= w) xo = xo - w; 
     for (int dy = -1; dy < 2; dy++) { 
     if (dx == 0 && dy == 0) continue; 
     cnn += cn; 
     yo = y + dy; 
     if (yo < 0) yo = h + yo; else if (yo >= h) yo = yo - h; 
     if (Mpart > 0) { Madd = 1; Mpart--;} else Madd = 0; 
     ii = (xo * h + yo) + cnn; 
     cells[ii] = make_int4(Fc, Mbase + Madd, Tc, 0); 
     } 
    } 
""", "ca_emit", preamble=""" 
#define Tk %s 
#define CN %s 
""" % (Tk, CELL_LEN)) 

absorb_gpu = ElementwiseKernel("int4 *cells, int *img, int w, int h, int cn", """ 
    int ii = i; 
    int4 cell = cells[i]; 
    int Fc = cell.x; 
    int Mc = cell.y; 
    int Tc = cell.z; 

    for (int c=0; c < 8; c++){ 
     ii += cn; 
     cell = cells[ii]; 
     int Fi = cell.x; 
     int Mi = cell.y; 
     int Ti = cell.z; 

     int dF = Fi - Fc; 
     if (dF > 180) Fc += 360; 
     if (dF < -180) Fc -= 360; 
     float sM = Mi + Mc; 
     if (sM != 0) sM = Mi/sM; else sM = 0; 
     dF = (int) (Fi - Fc) * sM; 
     int dM = Mi; 
     int dT = fabs((float) (Fi - Fc)) * fmin((float) Mc, (float) Mi)/Mk + (Ti - Tc) * sM; 
     Fc += dF; 
     Mc += dM; 
     Tc += dT; 
     Fc = Fc % 360; 
     if (Fc < 0) Fc += 360; 
     if (Tc > Tk) Tc = Tk; 
    }  

    cells[i] = make_int4(Fc, Mc, Tc, 0); 
    img[i] = hsv2rgb(Fc, Tc, Mc); 

""", "ca_absorb", preamble=""" 
#include <math.h> 
#define Mk %s 
#define Tk %s 

__device__ uint hsv2rgb(int hue, int sat, int val) { 
    // skipped for brevity 
} 
""" % (Mk, Tk))

감사 및 일부 최적화 문제에 대한 알렉세이 Shchepin에 : 은 [5 배 속도 향상은]

최적화 된 코드처럼 보인다.

출처

2014-12-16 23:43:56 a5kin

PyCuda를 이용한 유전자 세포 자동화, CUDA 커널에 셀당 많은 데이터를 효율적으로 전달하는 방법은 무엇입니까?

답변

관련 문제