Vote functions in a warp-divergent branch? Are they allowed? How idle threads are handled?

Accelerated Computing CUDA CUDA Programming and Performance
1

Quick question: are warp vote functions allowed in a branch with not all threads in a warp active? For example, how this is going to evaluate? :

if (threadIdx.x%2==0) {

  if (__any(some_predicate) ) {

	...

  }

  if (__all(some_other_predicate) ) {

	...

  }

}

In the programming guide I could only find that:

2

Any chance to get an answer?

3

I was wondering the same thing the other day and tried it out. Indeed, __any() and __all() can be used in divergent code. This is different from __syncthreads() which would deadlock in such a situation.

As to what the results will be: I have not tested this extensively but at least for __any(), the predicate seems to always evaluate to false for threads not participating in the execution. This seems logical to me and was exactly what I was looking for.

Since it appears safe to call __any() and __all() from divergent code, you can do your own experiments to see what results they produce.

4

The definition from the programming guide is ambiguous, but the PTX manual has a more accurate description:

vote		Vote across thread group.

Description Performs a reduction of the source predicate across threads in a warp. The destination

			predicate value is the same across all threads in the warp.

			The reduction modes are:

			.all True if source predicate is True for all active threads in warp. Negate the source

				  predicate to compute .none.

			.any True if source predicate is True for some active thread in warp. Negate the

				  source predicate to compute .not_all.

			.uni True if source predicate has the same value in all active threads in warp.

				  Negating the source predicate also computes .uni.

Note the “active” word…

5

nVidia has some example code at
[url=“http://www.nvidia.com/content/cudazone/cuda_sdk/CUDA_Basic_Topics.html#simpleVoteIntrinsics”]http://www.nvidia.com/content/cudazone/cud...eVoteIntrinsics[/url]
Bill

    Dr. W. B. Langdon, 
    Department of Computer Science,
    University College London
    Gower Street, London WC1E 6BT, UK
    [url="http://www.cs.ucl.ac.uk/staff/W.Langdon/"]http://www.cs.ucl.ac.uk/staff/W.Langdon/[/url]
6

nVidia has some example code at
[url=“http://www.nvidia.com/content/cudazone/cuda_sdk/CUDA_Basic_Topics.html#simpleVoteIntrinsics”]http://www.nvidia.com/content/cudazone/cud...eVoteIntrinsics[/url]
Bill

    Dr. W. B. Langdon, 
    Department of Computer Science,
    University College London
    Gower Street, London WC1E 6BT, UK
    [url="http://www.cs.ucl.ac.uk/staff/W.Langdon/"]http://www.cs.ucl.ac.uk/staff/W.Langdon/[/url]