19903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski// RUN: %clang_cc1 -triple nvptx-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s 29903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s 39903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski 49903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinskiint bar(int a) { 59903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski int result; 69903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski // CHECK: call i32 asm sideeffect "{ {{.*}} 79903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski asm __volatile__ ("{ \n\t" 89903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski ".reg .pred \t%%p1; \n\t" 99903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski ".reg .pred \t%%p2; \n\t" 109903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski "setp.ne.u32 \t%%p1, %1, 0; \n\t" 119903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski "vote.any.pred \t%%p2, %%p1; \n\t" 129903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski "selp.s32 \t%0, 1, 0, %%p2; \n\t" 139903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski "}" : "=r"(result) : "r"(a)); 149903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski return result; 159903e94bee8fbbec6d44fe2de91ccee46f40187fJustin Holewinski} 16