1*f4a2713aSLionel Sambuc // RUN: %clang_cc1 -triple nvptx-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s 2*f4a2713aSLionel Sambuc // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -O3 -S -o - %s -emit-llvm | FileCheck %s 3*f4a2713aSLionel Sambuc bar(int a)4*f4a2713aSLionel Sambucint bar(int a) { 5*f4a2713aSLionel Sambuc int result; 6*f4a2713aSLionel Sambuc // CHECK: call i32 asm sideeffect "{ {{.*}} 7*f4a2713aSLionel Sambuc asm __volatile__ ("{ \n\t" 8*f4a2713aSLionel Sambuc ".reg .pred \t%%p1; \n\t" 9*f4a2713aSLionel Sambuc ".reg .pred \t%%p2; \n\t" 10*f4a2713aSLionel Sambuc "setp.ne.u32 \t%%p1, %1, 0; \n\t" 11*f4a2713aSLionel Sambuc "vote.any.pred \t%%p2, %%p1; \n\t" 12*f4a2713aSLionel Sambuc "selp.s32 \t%0, 1, 0, %%p2; \n\t" 13*f4a2713aSLionel Sambuc "}" : "=r"(result) : "r"(a)); 14*f4a2713aSLionel Sambuc return result; 15*f4a2713aSLionel Sambuc } 16