20 replies to this topic

[Alienizer]

I have a very simple kernel that is puzzling me all night long. I Pass 2 write arrays of doubl4, and one read array double4 of random numbers 0..1. I also pass a (double) angle value in rad. The problem is that one specific line works, the other does not. What I mean by that is, I'm returning the array of random numbers, but they are all zero. With the other line, it return the correct random value. Please see the code, I marke the 2 lines of wich work and not...

__kernel void testing(
			__global double4* dstA,	// host set as CL_MEM_WRITE_ONLY
			__global double4* dstB,	// host set as CL_MEM_WRITE_ONLY
			__global const double4* rnd,  // host set as CL_MEM_READ_ONLY
			int iNumElements,
			double angleInRad
		   )
{
int iGID = get_global_id(0);
if (iGID >= iNumElements){
  return;
}
double4 rnd = rnd[iGID];
//double4 org = (double4)(0,0,0,0);  // this works
double4 org = (double4)(sin(angleInRad), 0,0,0); // this doesn't work
//double4 org = (double4)(sqrt(angleInRad), 0,0,0); // but this work
dstA[iGID] = org;
dstB[iGID] = rnd;
}

If I enable this line and rem out the line after...
double4 org = (double4)(0,0,0,0);
//double4 org = (double4)(sin(angleInRad), 0,0,0);
the kernel return the original array of random mnumbers. But if I do the inverse, the array returned is all zero!!??

Any ideas as to what I'm doing wrong because I can't see how the sin() function (which works) is causing everything to be zero!

[Vilem Otte]

So, I tried the code on my notebook (as I'm not at home right now) and it seems there are few issues in it.

First one (that I had to solve) was mine, I don't have doubles supported on GPU, so I had to use floats. Another one is that passing __global float4* dstA couldn't be accessed on my machine like dstA[i] -> so I passed it through my own structure (vec4 - defined as 'typedef struct { float x, y, z, w; } vec4;'). Anyway I hope these minor changes didn't affect the result (or the problem). Anyway...

The problem was (and it didn't even compile on AMD OpenCL compiler), that you're passing __global const double4* rnd; and you're also declaring double4 rnd = rnd[iGID]; - e.g. you're naming two variables with the same name (rnd) - C99 doesn't allow it, neither OpenCL specs. I don't know why it even compiled on NVidia card (assuming you have one, based on previous thread) - as it should give you at least warring, better error. Try to rename it and see what it does.

[Reedbeta]

Also, a note about doubles - GPUs are mainly designed for single-precision and even half-precision floats; some don't even support doubles (as Vilem found) and even when they do, double-precision math is likely to be a lot slower than float or half. They consume twice the memory bandwidth too, which doesn't help anything. So, try to avoid using doubles unless it's really necessary for the extra precision and range.

[Alienizer]

Thank you both for the help. I changed the code to use only float, no double (the rnd syntax was in the post only, notin the original code). But here is the original code now, but the sincos() function makes it crash! any of the sin() cos() sincos() functions make it crash. If I use a sqrt instead, the code works!?

 
typedef struct { float x, y, z, w; } vec4;

__kernel void Testing(
			__global vec4* dst,
			__global const vec4* rnd,
			int iNumElements
		   )
{
 int iGID = get_global_id(0);

 if (iGID >= iNumElements){
  return;
 }

  vec4 rndVec = rnd[iGID];

  float r = rndVec.x;
  float q = rndVec.y;

  float c;
  float s = sincos(q, &c);

  vec4 org = (vec4){ r, q, 0, 0 };
  dst[iGID] = org;
}

[Vilem Otte]

Hm... this definitely look very strange (as the same code works on OpenCL 1.0 and 1.1 (haven't tested on 1.2 as it is just in beta App SDK, and we need stable release in work (also the 1.2 supports C++98 standard, and we're developing modules that are working with OpenCL in Gnu99 (C99) standard)). This seem as driver-related issue.

I went through NVidia forums and it seem that sin/cos functions are not supported for double-precision (but it was from March 2011 - so I don't know whether they support it now), but their OpenCL reference specifies that sin/cos functions are defined and specified for single-precision floats - so your code should work (as it works on AMD gpus).
If you haven't done already, try updating drivers, but I presume you're up-to-date - so I'd try contacting NVidia and in time between I'd write my own sin/cos implementation (it is just some multiply-add operations), sin for full period is defined as:

So it's implemetation is actually quite simple (and you can pre-compute those factorials (or better 1.0/3!, 1.0/5! and 1.0/7! - as you won't need divisions then, just multiplies) to speed computation a lot). I think that GPU even does the Taylor expansion internally (as neither sin or cos are hardware computed).

[Alienizer]

It's really weird! Even the following super simple code crash! But if I change the sin function for a sqrt function, then it doesn't crash.

And what's more weird, if I put in say 0.5 in the sin() function as oppose to angleInRad, then it works! I've even tested that angleInRad = 0.5 to make sure.

__kernel void Testing( float angleInRad )
{

  float s = sin(angleInRad);

}

Does that code run on yours?

[Vilem Otte]

Yup, this simple testing kernel runs (without any warning or error). This seems weird. Could you post how are you passing/reading your variables to OpenCL application (just to make sure there isn't any mistake ... though if sqrt is working and sin not then it shouldn't be issue)?

[Alienizer]

I know, it's weird! It does the same on my other PC with an ATI. I use the NVIDIA GPU Computing SDK, oclVectorAdd sample and modified it, the sample itself runs, but not when there is a sin() with a var as a parameter, only works with const.

float angleInRad = 0.5;
clSetKernelArg(ckKernel, 0, sizeof(cl_float), (void*)&angleInRad);

[TheNut]

I think your crashing is memory related. I see you're using a 4-vector float, which means you should have allocated [SIZE * 4 * sizeof(float)] in your device and host memory. Since you're only using the first index in the vector while leaving the other indices zeroed, I'm willing to believe you didn't supply a sufficiently large enough memory block. When the global_id passes [SIZE / 4], it runs out of memory and screams bloody murder Do you mind posting your C++ OpenCL code? It could shed some clues.

Also, for efficiency reasons you should maximize the use of vectors instead of computing the trig on individual floats. For example, let's say you want to calculate 1024 sin and cos values. This is how I would set it up.

In C++,

// Allocate host memory
unsigned int size = 1024;
float *hostRead = new float[size];
float *hostWriteSin = new float[size];
float *hostWriteCos = new float[size];

// Initialize the read array
for (unsigned int i = 0; i < size; ++i)
{
	 hostRead[i] = (i / (float)size) * Math::PI;
}

// Create device memory and copy values from host memory
cl_mem memRead= clCreateBuffer(clContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, size * sizeof(float), hostRead, &result);
cl_mem memWriteSin = clCreateBuffer(clContext, CL_MEM_WRITE_ONLY, size * sizeof(float), 0, &result);
cl_mem memWriteCos = clCreateBuffer(clContext, CL_MEM_WRITE_ONLY, size * sizeof(float), 0, &result);

// Run kernel, copy back device memory to host memory, etc..

You know you have 1024 values, but you can reduce your calculations to 1024 / 4 = 256 (due to simd architecture). So, your kernel should look like this:

__kernel void vector_gpu (__global const float4 *src,
		 __global float4 *dst,
		 __global float4 *dst2,
		 const int num)
{
	 // Where "num" == 1024 / 4 = 256. Only 256 executions.
	 const int globalId = get_global_id(0);
	 if ( globalId < num )
		  dst[globalId] = sincos(src[globalId], &dst2[globalId]);
}

[Vilem Otte]

Just a note...

TheNut, on 08 January 2012 - 05:21 PM, said:

I'm surprised your kernel compiles, since vec4 is not part of the standard.

I posted few posts ago, that he can use typedef struct { float x, y, z, w; } vec4; instead of builtin float4 (I think some older APP sdk don't allow float4 - so I thought it might be the same issue). Also it will be compiled to vector instructions, so it won't give you speed penalty.

[TheNut]

I noticed that after I made the post. It just escaped my vision when I scanned through his code. I recently edited out that part in my reply.

[Reedbeta]

Just a note, the reason that sin(aConstant) fails when sin(aVariable) works is likely that the compiler is performing constant-folding optimization and precomputing all the values that depend on constants. So the kernel ends up not actually evaluating sin at all. As for why sqrt works and sin doesn't, that indeed sounds like a driver bug.

[Alienizer]

Reedbeta, on 08 January 2012 - 06:23 PM, said:

Just a note, the reason that sin(aConstant) fails when sin(aVariable) works is likely that the compiler is performing constant-folding optimization and precomputing all the values that depend on constants. So the kernel ends up not actually evaluating sin at all. As for why sqrt works and sin doesn't, that indeed sounds like a driver bug.

Actualy, it works with constant, but doesn't with var!

[Reedbeta]

Yes, that's what I meant. I said it backward, but my explanation was correct.

[Alienizer]

Reedbeta, on 08 January 2012 - 06:23 PM, said:

As for why sqrt works and sin doesn't, that indeed sounds like a driver bug.

On all my 3 PCs??? each have differnt video cards, 2 Nvidia and 1 ATI

Does it work for you?

[Stainless]

I have had shader compiler bugs in the past.

When I was working on my atmospheric scattering shader in HLSL I had a bug that took me ages to find which I finally tracked down to a single line of code.

I had float4 calculated value * = float and I changed it to float4 var = calculated value * float which fixed the problem. Note I use *= all over the place and it works, it was just the combination of code before and after that LOC which screwed the compiler.

So what I would advise you to do is back up your code, then rip everything out.

Start with the simplest bit of code you can. No typedefs, no subroutines, just define shader, return sin of something.

If that works start adding code a tiny bit at a time until it breaks.

I know it's a pain in the proverbial, but without a proper shader debugger it's the only way to find problems like this.

[Vilem Otte]

#Stainless - Actually there is a debugger for OpenCL - http://developer.amd...es/ocl-emu.aspx

As for shaders, I thought of designing and writing proper shader debugger once or twice (even came with some kind of design that would work) but well I didn't actually have time to build it (nor time for trying to search for possible co-workers).

[Alienizer]

I know what you mean Stainless, but the code I used is a very very simple VectorAdd sample from Nvidia SDK, which compiles and run perfect. But if I change the add function for a sin function, then it no longer works, but it does with sqrt. In all my PCs, ATI and Nvidia, with all updated drivers. It's driving me so crazy you have no idea! I'll try the debugger Vilem Otte suggested.

[TheNut]

Alienizer said:

It's driving me so crazy you have no idea

We have a pretty good idea Happens to all of us once in a while. Have you tried polling the program log after you compile and build the program? It should give you a pretty detailed list of compilation warnings that might help. If you try out the debugger, try compiling and linking your program against the OpenCL library in the AMD APP SDK and see if that helps. I have no problems running all kinds of math routines on my AMD card, including sin and sincos.

[Alienizer]

I just figure part of it out!

The crash is when calling clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);

If I put in '-cl-fast-relaxed-math' as the options, then it doesn't crash at that point, only later, not sure yet.

So it seems I need some options, the right one(s) perhaps??