numeric indexing vector component

Accelerated Computing CUDA CUDA Programming and Performance
1

Are you allowed to index a vector’s components with operator as if they were members of an array? You can in D3D10/11 shader. With my nvidia OpenCL device I try it, it compiles, but it gives wrong results. eg:

uint4 a, b;



/*

#pragma unroll

for(int i = 0; i < 4; ++i) {

	a[i] = pass1_values;

	b[i] = pass1_values;

}*/



a.x = pass1_values;

a.y = pass1_values;

a.z = pass1_values;

a.w = pass1_values;

b.x = pass1_values;

b.y = pass1_values;

b.z = pass1_values;

b.w = pass1_values;

The commented portion should do the same (I think) as the uncommented portion, but gives apparently undefined results. It compiles without warning.

.sean

2

Hi Sean,
I have tried your example (replacing pass1_values by simple values), I get correct results both on uint4 and float4.
Perhaps you can give more details.

I’m running GPU computing SDK 3.0 on driver 195.81, win7 64, AMD CPU and 275 GTX GPU.

Have to admit that I had to correct some errors I made in the code calling the kernel and retrieving/displaying the results before it worked out correctly… External Image

3

Hmm. It didn’t work for me. I also tried on my ATI 5800, and same thing: compiled but didn’t work. I posted to the ATI board and got this response:
[url=“http://forums.amd.com/devforum/messageview.cfm?catid=390&threadid=124303&enterthread=y”]http://forums.amd.com/devforum/messageview...p;enterthread=y[/url]

ATI considers it illegal syntax and will update the compiler to give an error in the next release. It is very frustrating that there is no reference implementation like D3D has.

4

Thanks for your post, saves me probably from code being broken at some future revision.

I could give you my code, of course, just out of curiosity. Afa Micah’s response, that was at first a bit baffling to me.

Here’s an excerpt from OpenCL spec 6.1.7:

[i]Elements of vector data types can also be accessed using a numeric index to refer to the

appropriate element in the vector. The numeric indices that can be used are given in the table

below:

Vector Components Numeric indices that can be used

2-component 0, 1

4-component 0, 1, 2, 3

8-component 0, 1, 2, 3, 4, 5, 6, 7

16-component 0, 1, 2, 3, 4, 5, 6, 7,8, 9, a, A, b, B, c, C, d, D, e, E,f, F

The numeric indices must be preceded by the letter s or S.[/i]

Only trouble is that appears to be illegal, heavens knows why. Instead of a[0] a.s0 is to be used.

So it is interesting to see if an int variable can be used as index.

This causes an error, as was to be expected:

:187: error: illegal vector component name 'i'

				fl.si = anumber + i;

				  ^~~

So it looks as if a loop cannot be used to sequentially assign components of a vector (without explicitily naming them).

You could still use an ordinary array of ints instead of a vector-type. Perhaps that makes no fundamental difference for your code?

You would declare unsigned int a[4], b[4]; ?

But I agree that it seems a frustrating omission if indexing vectorelements by means of a variable is not implemented. Maybe there is a good reason, but I don’t know it.

Jan

5

I don’t know if this makes you any happier, since it involves a pointer-cast, but the following actually works:

__kernel void test	(

				__global uint4* uintg

				)

{

	uint u[4];

	#pragma unroll

	for(int i = 0; i < 4; ++i)

	{

		u[i]= i+1;

	}

	*uintg=*(uint4*)u;

}

So, if you are willing to use this, you could try:

a = *(uint4*)&pass1_values;

assuming pass1_values is an array of uints.

You may have to pay attention to the memory space, e.g. (__global uint4) if your array is in global memory.

Jan