PTX asm constraint letter for predicate

Accelerated Computing CUDA CUDA Programming and Performance
1

What is PTX asm constraint letter for predicate?

I have the following code

__global__ void f(){
    bool p, q;
    asm volatile("vote.sync.all.pred %0, %1, 0xffffffff;" : "=r"(p) : "r"(q));
}

This don’t work because error: asm operand type size(1) does not match type/size implied by constraint 'r'

But I can’t find the corresponding constraint letter, all that’s listed is here

"h" = .u16 reg
"r" = .u32 reg
"l" = .u64 reg
"f" = .f32 reg
"d" = .f64 reg
2

As njuffa showed at this recent discussion, now nvidia doesn’t support defining predicate registers in a single-line inline ptx assembly, you can use multi-line assembly and convert it to a 32-bit register type with the selp instruction.
Here is the code I adapted from his example, I have confirmed that it will compile successfully, but there is no verification of correctness.

__global__ void f(){
	unsigned p, q;
	asm ("\n\t"
		 ".reg .pred p, q;\n\t"
		 "setp.ne.b32  q, %1, 0;\n\t"
		 "vote.sync.all.pred p, q, 0xffffffff;\n\t"
		 "selp.u32     %0, 1, 0, p;\n\t"
		 : "=r"(p) : "r"(q));
}
3

Thanks! So writing ptx is not really convenient for some cases I guess.

4

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.