배정 밀도 데이터를위한 CUDA의 워프 셔플는

CUDA는 배정 밀도 데이터를 감소해야 프로그램, 나는 줄리앙 Demouth의 이름 슬라이드 사용 "셔플 : 팁과 트릭"배정 밀도 데이터를위한 CUDA의 워프 셔플는

을 셔플 기능은 다음과 같습니다 :

/*for shuffle of double-precision point */ 
__device__ __inline__ double shfl(double x, int lane) 
{ 
    int warpSize = 32; 
    // Split the double number into 2 32b registers. 
    int lo, hi; 
    asm volatile("mov.b32 {%0,%1}, %2;":"=r"(lo),"=r"(hi):"d"(x)); 
    // Shuffle the two 32b registers. 
    lo = __shfl_xor(lo,lane,warpSize); 
    hi = __shfl_xor(hi,lane,warpSize); 
    // Recreate the 64b number. 
    asm volatile("mov.b64 %0,{%1,%2};":"=d"(x):"r"(lo),"r"(hi)); 
    return x; 
}

현재 프로그램을 컴파일하는 동안 아래 오류가 있습니다.

ptxas /tmp/tmpxft_00002cfb_00000000-5_csr_double.ptx, line 71; error : Arguments mismatch for instruction 'mov' 
ptxas /tmp/tmpxft_00002cfb_00000000-5_csr_double.ptx, line 271; error : Arguments mismatch for instruction 'mov' 
ptxas /tmp/tmpxft_00002cfb_00000000-5_csr_double.ptx, line 287; error : Arguments mismatch for instruction 'mov' 
ptxas /tmp/tmpxft_00002cfb_00000000-5_csr_double.ptx, line 302; error : Arguments mismatch for instruction 'mov' 
ptxas /tmp/tmpxft_00002cfb_00000000-5_csr_double.ptx, line 317; error : Arguments mismatch for instruction 'mov' 
ptxas /tmp/tmpxft_00002cfb_00000000-5_csr_double.ptx, line 332; error : Arguments mismatch for instruction 'mov' 
ptxas fatal : Ptx assembly aborted due to errors 
make: *** [csr_double] error 255

누군가가 조언을 줄 수 있습니까?

출처

2014-06-07 taoyuan

귀하의 문제는 중괄호라고 생각합니다. 나는 그들이 gcc (최소한 i386에서)에 특정한 의미를 가지고 있으며 어셈블러로 전달되기 전에 소비된다는 것을 믿는다. -S 출력이 어떻게 보이는지 보는 것은 흥미로울 것입니다. 특히 인라인 asm에 대한 APP 섹션. –

@DavidWohlferd : 중괄호는 완전하며 코드를 컴파일하는 gcc가 아닙니다. NVIDIA의 lvmm 기반 GPU 컴파일러입니다. – talonmies

@talonmies 더 의미가 있습니다. 확실히 gcc처럼 보일지 모르지만, 그렇다면 절대로 작동하지 않을 수 있습니다. –

32 비트 레지스터에 대한 double 인수의로드에 대한 인라인 어셈블리 명령에 구문 오류가 있습니다. 이것 (즉 64 비트 부동 소수점 레지스터) 32 비트 부하 소스로 불법적 인 "D"를 사용

asm volatile("mov.b64 {%0,%1}, %2;":"=r"(lo),"=r"(hi):"d"(x));

(및 mov.b32 아니오한다 :

asm volatile("mov.b32 {%0,%1}, %2;":"=r"(lo),"=r"(hi):"d"(x));

같아야 여기에서 코드는 두 개의 32 비트 레지스터에 64 비트를로드해야 함).

출처

2014-06-07 08:52:23 talonmies

정말 고마워! – taoyuan

@ user3201449 : 또한이 함수의 지역 변수'warpSize'가 내장 변수'warpSize'의 별칭임을 지적 할 가치가 있습니다. 그 줄은 제거해야합니다. 그것은 오늘날 어떤 것도 깨뜨리지 않을 것이지만, 언젠가 미래 아키텍처가 어떻게 될지에 따라 달라질 것입니다. – talonmies

@tolonmies : 귀하의 지시에 따라 삭제했습니다. 어쨌든 감사합니다! – taoyuan

CUDA 9.0부터, __shfl, __shfl_up, __shfl_down 및 __shfl_xor은 더 이상 사용되지 않습니다.

새로 도입 된 기능 __shfl_sync, __shfl_up_sync, __shfl_down_sync 및 __shfl_xor_sync 다음과 같은 프로토 타입을 가지고 T이 int, unsigned int, long, unsigned long, long long, unsigned long long, float 또는 double 수 있습니다

T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize); 
T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize); 
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int 
width=warpSize); 
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);

더 이상 배정도 산술을위한 임의의 셔플 명령어를 만들 필요가 없습니다.

출처

2018-03-08 07:37:29 JackOLantern

배정 밀도 데이터를위한 CUDA의 워프 셔플는

답변

관련 문제