cuda fortran의 공유 메모리가 예상대로 작동하지 않습니다.

Question

나는 cuda fortran을 구축 중이며 이상한 동작이 발생합니다. 나는 왜 내 코드가 이런 식으로 돌아가는지 이해하지 못하고 당신의 도움에 감사 할 것입니다.

0 값이 지정되지 않고 심지어 루프 이 경계를 넘어서 실행됩니다.

루프 뒤에 if 조건을 넣으려고했으나 도움이되지 않았습니다. 내가 할 모든 글로벌 메모리 어레이의 크기를 설정

real, shared :: s_d_aaa_adk(0:15,0:15) 
    real, shared :: s_d_bbb_adk(0:15,0:15) 
    real, shared :: s_d_ccc_adk(0:15,0:15) 

    d_k = (blockIdx%x-1) 
    s_d_j = threadIdx%x-1 
    s_d_l = threadIdx%y-1 

    if(d_k == kmax-1)then 
     s_d_aaa_adk(s_d_j,s_d_l) = 0 
     s_d_bbb_adk(s_d_j,s_d_l) = 0 
     s_d_ccc_adk(s_d_j,s_d_l) = 0  
    endif 

    do d_k = 0, kmax-2    
     s_d_bbb_adk(s_d_j,s_d_l) = d_bbb(s_d_j,d_l,d_k+1) 
     s_d_ccc_adk(s_d_j,s_d_l) = d_ccc(d_j,s_d_l,d_k+1) 
     s_d_aaa_adk(s_d_j,s_d_l) = d_aaa(d_j,s_d_l,d_k+1)    
    end do ` 

당신의 도움 (16, 16,의 Kmax), 그리드 (128,1,1)는, 블록 (16,16,1)을 주셔서 감사합니다 그리고 커널은 블록 인덱스에서 파생 당신이있는 거 조절은 if 문 d_k에, 이후 testkernell<<<grid,block>>>()

Answer 1

로 시작되는이 의미

d_k = (blockIdx%x-1) 
if(d_k == kmax-1)then 

그 (128) 중 하나의 블록 당신의 그리드에서 실제로 if 문을 실행하여 해당 공유 메모리 값을 0으로 설정합니다. 대부분의 블록은 if 문 안에있는 것을 실행하지 않습니다.

그리고 kmax이 128보다 큰 경우, 블록 중 어느 것도 if 문을 실행하지 않습니다.

if 문을 모든 스레드 블록 내에서 실행하려면 블록 인덱스가 아닌 다른 명령문에서 조건문을 처리해야합니다.

코드를 재구성하는 방법에 대한 제안을했지만, 공유 메모리에 데이터를로드하는 것까지 달성하고자하는 것이 분명하지 않습니다. 예를 들어, DO 루프 나에게 많은 이해가되지 않습니다 :

do d_k = 0, kmax-2    
    s_d_bbb_adk(s_d_j,s_d_l) = d_bbb(s_d_j,d_l,d_k+1) 
    s_d_ccc_adk(s_d_j,s_d_l) = d_ccc(d_j,s_d_l,d_k+1) 
    s_d_aaa_adk(s_d_j,s_d_l) = d_aaa(d_j,s_d_l,d_k+1)    
end do   ^ ^
        |  | 
     a given thread has specific values for these indices 

귀하의 s_d_j 및 s_d_l 변수는 스레드 인덱스입니다. 그래서 주어진 스레드는이 do 루프를 보게 될 것이고 루프를 반복적으로 실행하여 다양한 전역 메모리 배열 (d_bbb, d_ccc 등)의 연속적인 값을 정확히 일치하는 각 공유 메모리 배열의 위치에로드합니다.

스레드 실행이 실제로 어떻게 작동하는지 이해하지 못하는 것 같습니다. 당신이 주어진 스레드라고 가정하고 s_d_j 및 s_d_l (및 d_k)에 특정 값을 지정하십시오. 루프 인덱스로 해당 변수를 재사용 할 때 블록 인덱스를 덮어 쓰는 경우도 있지만 이상하게 보입니다. 코드 실행은 의미가 있습니다.

편집 : 추가 의견을 바탕으로 :

당신은 당신의 전체 데이터 세트의 크기를 (x는, Y, Z)이 (64,64,32)는 언급했다. "각 슬라이스를 하나의 블록에 넣고 싶습니다."

슬라이스 당 하나의 블록을 시작해야한다고 제안했습니다. 또는 단일 슬라이스에 여러 블록이 할당 된 알고리즘을 염두에 두어야합니다. 그럼에도 불구하고, 나는 그 슬라이스에 할당 된 주어진 블록에 모든 슬라이스 데이터 (64, 64)를 사용할 수 있기를 원한다고 가정합니다. 나는 지금 당신이 32 블록을 시작할 것으로 가정 할 것이다.단일 슬라이스에서 여러 블록이 작동하는 경우로 확장하는 것이 어렵지 않아야합니다. 나는 또한 지시 한 16x16보다 32x32 스레드 블록을 가정 할 것입니다. 원하는 경우 16x16을 사용하도록 확장하는 것이 어렵지 않습니다.

real, shared :: s_d_aaa_adk(0:63,0:63) 
real, shared :: s_d_bbb_adk(0:63,0:63) 
real, shared :: s_d_ccc_adk(0:63,0:63) 

c above uses 48KB of shared mem, so assuming cc 2.0+ and cache config set accordingly 

d_k = (blockIdx%x-1) 
s_d_j = threadIdx%x-1 
s_d_l = threadIdx%y-1 

c fill first quadrant 
s_d_bbb_adk(s_d_j,s_d_l) = d_bbb(s_d_j,s_d_l,d_k+1) 
s_d_ccc_adk(s_d_j,s_d_l) = d_ccc(s_d_j,s_d_l,d_k+1) 
s_d_aaa_adk(s_d_j,s_d_l) = d_aaa(s_d_j,s_d_l,d_k+1) 
c fill second quadrant 
s_d_bbb_adk(s_d_j+blockDim%x,s_d_l) = d_bbb(s_d_j+blockDim%x,s_d_l,d_k+1) 
s_d_ccc_adk(s_d_j+blockDim%x,s_d_l) = d_ccc(s_d_j+blockDim%x,s_d_l,d_k+1) 
s_d_aaa_adk(s_d_j+blockDim%x,s_d_l) = d_aaa(s_d_j+blockDim%x,s_d_l,d_k+1) 
c fill third quadrant 
s_d_bbb_adk(s_d_j,s_d_l+blockDim%y) = d_bbb(s_d_j,s_d_l+blockDim%y,d_k+1) 
s_d_ccc_adk(s_d_j,s_d_l+blockDim%y) = d_ccc(s_d_j,s_d_l+blockDim%y,d_k+1) 
s_d_aaa_adk(s_d_j,s_d_l+blockDim%y) = d_aaa(s_d_j,s_d_l+blockDim%y,d_k+1) 
c fill fourth quadrant 
s_d_bbb_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = d_bbb(s_d_j+blockDim%x,s_d_l+blockDim%y,d_k+1) 
s_d_ccc_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = d_ccc(s_d_j+blockDim%x,s_d_l+blockDim%y,d_k+1) 
s_d_aaa_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = d_aaa(s_d_j+blockDim%x,s_d_l+blockDim%y,d_k+1) 


c just guessing about what your intent was on filling with zeroes 
c this just makes sure that one of the slices at the end gets zeroes 
c instead of the values from the global arrays 

if(d_k == kmax-1)then 
c fill first quadrant 
    s_d_bbb_adk(s_d_j,s_d_l) = 0 
    s_d_ccc_adk(s_d_j,s_d_l) = 0 
    s_d_aaa_adk(s_d_j,s_d_l) = 0 
c fill second quadrant 
    s_d_bbb_adk(s_d_j+blockDim%x,s_d_l) = 0 
    s_d_ccc_adk(s_d_j+blockDim%x,s_d_l) = 0 
    s_d_aaa_adk(s_d_j+blockDim%x,s_d_l) = 0 
c fill third quadrant 
    s_d_bbb_adk(s_d_j,s_d_l+blockDim%y) = 0 
    s_d_ccc_adk(s_d_j,s_d_l+blockDim%y) = 0 
    s_d_aaa_adk(s_d_j,s_d_l+blockDim%y) = 0 
c fill fourth quadrant 
    s_d_bbb_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = 0 
    s_d_ccc_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = 0 
    s_d_aaa_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = 0  
endif