I would like to to use operator && two int4 and expect a int4 as result. I get an error, please see attached message.

Have you any hints?

Thanks

screw

I am a student working on a research project and I right now I have to write a model reduced fluid simulation on the GPU. Exactly what must be done on the GPU follows:

(1) look up the corresponding tet via some kind of spatial data structure
(2) do a matrix multiply with the reduced state to figure out the fluxes
(3) convert the fluxes to a velocity via a 3x3 matrix inversion
(4) optionally: cache the velocity for this timestep
(5) advect the particle
(6) resolve any collisions with the fixed geometry (car)
(7) render all the particles and a car

do you think i would get big advanteges from the use of OpenCL respect to DirectX11?

What about the rendering?

thanks

I have a machine with 4GB RAM memory and HD5870 with 1GB memory. When I use clGetDeviceInfo to get CL_DEVICE_GLOBAL_MEM_SIZE it only return 3GB RAM in my machine and only 256MB video memory from HD5870 which is actually much smaller than the original value.

At first I though memory used by OS will deducted but end up the value getting was not.

When I try to allocate 140MB memory from HD5870, it fail and said memory not enough. In the end, I was able to use less than 110MB than 1GB memory provided by the graphic card.

I wish someone can tell me what is actually happen, that will be gratefully to help me understanding more in OpenCL.

Thank you.

Should we redownload them or has nothing inside these files been changed?

Also does the Catalyst 9.11 driver include OpenCL support?

However, I am testing on 2 machines that have NVIDIA cards and Intel CPUs, and the samples run fine with DEVICE_TYPE_GPU, which impresses me, but with DEVICE_TYPE_CPU the call to clCreateContextFromType()  fails with CL_DEVICE_NOT_FOUND.

On one machine I went ahead and installed the driver package (Catalyst), but that made no difference.

The machines are Vista64 and the builds are 64 bit.

What is it that I don't get? 
Thanks in advance!

What i would want to know from community experts is that, how will i be able to port a part of a complex computation to the GPU? Like, say if i'm doin a triple integration, with all limits from -infinity to +infinity, will i be able to port two inner integrals to the GPU and then then last one to the CPU?

I also wanted to learn about how to use the Stream SDK. Any help or links would be most welcome. I'l be really thankful

Regards,

Ankur

The results I get for my system (cpu=Dual Core Pentium E5200, gpu=HD4770):

Local memory:
cpu: 32768
gpu: 16384
Memory type:
cpu: 2
gpu: 2
         cpu native - exec.time:       0.24 t2:          0          1          2          3
           cpu copy - exec.time:       0.07 t2:          0          1          2          3
         cpu copy4 - exec.time:       0.03 t2:          0          1          2          3
cpu copy4_local - exec.time:       0.03 t2:          0          1          2          3
           gpu copy - exec.time:       49.3 t2:          0          1          2          3
         gpu copy4 - exec.time:       15.6 t2:          0          1          2          3
gpu copy4_local - exec.time:       15.8 t2:          0          1          2          3


Conclusions:
1) OpenCL implementation substantially outperformed native implementation,
about 3x faster. Since the cpu has 2 cores, I'm not sure where this speed up
comes from.
2) Vectorized version performs better, as expected. About ~2.5-3x better
than scalar version.
3) gpu performance is not comparable to cpu. cpu is a few hundert times
faster.


I'm posting this results to hear experiences from other users. Furthermore,
it would be nice to see results of the same program on other systems,
particularly those which have dedicated(fast) local memory, which is not the case with my system.
Any comments on how to increase gpu performance with respect to cpu
performance are welcomed. I would also be grateful to those who could give me good explanation of the posted results.It would be nice to know why is cpu __local associated to CL_GLOBAL (explanation for GPU can be found in other threads).

P.S. Structure of the kernel is chosen so deliberately. It reflects more
complex structure of kernels I use. Of course, this simplified structure can
further be simplified, but then it wouldn't reflect demands imposed on
hardware by more complex kernels. Performance gains through async copying don't count either.)

If you want full OpenCL unfortunatelly it's true. But not every program needs full OpenCL. Usually only small subset of the language is required ( like for matrix multiplication, vector addition , etc ).

So lets consider if RV6xx can support some part of OpenCL.

RV6xx doesn't have compute shaders. It means that we can't give any NDRange ( work items index space ). Our NDRange is defined by output buffer size.

1. So we have first restriction here - NDRange == dimension & size of output buffer. This is usually the case when we do classical stream computing ( convert one stream into another - as in brook ).

2. Second limitation is that we can't request any local work group size - the driver makes decision here.

3. Reads from memory must use texture unit. We can't write to read buffers ( it implies that kernel must have const pointers to read buffers ).

4. Writes to output buffer must follow pattern

out[gloval_id(0)] = value; ( for 1D )

out[global_id(1)*out_width + global_id(0)] = value; ( for 2D )

5. We can't use local memory ( missing on RV6xx ).

So on device side we have kernels which are restricted to the form:

void kernel( const global floatx* input1 [, const global floatx* inputx], global floatx* output1 [, global floatx* outputx], floatx,intx variables ( not pointers ) )

{

    // can't use local_id, local_size

    // only global_id, global size available

   ... computations and memory read buffer access here ...

   output1[ global_id(1)*output_width + global_id(0)] = value; // 2D case

   // up to 8 outputs

}

such a kernel can be compilled into pixel shader. Kernels not matching this pattern should give compilation error on RV6xx.

On the host side we have limitation that NDRange given to kernel invocation must be as the same as output buffer size. ( there is some technical problem of opencl buffer not having 2D size, but it can be easilly solved by small extension to cal ).

So as we see some part of OpenCL can be supported on RV6xx. Now you can decide for yourself whether this model is sufficient for your work or not.

Personally i think that AMD/ATI should implement this. It is logical extension of Brook framework ( connecting old with new ). Also it would be nice gesture from AMD/ATI towards people with older cards.

Can anyone give me some pointer for which -if any- performance profiling tools work with  GPU OpenCL kernels? A measure of SP occupancy and access to any profiling counters like CUDA permits would do the job nicely.

Ta,

MJH

&::clGetKernelWorkGroupInfo, device(), object_, name, param),

should be replaced by

&::clGetKernelWorkGroupInfo, object_, device(), name, param).

Unfortunately, it seems that the kernels don't run concurrently, or else I'm timing something incorrectly.  What I'd like to see is the following (for instance):

Time for a single kernel to run on the GPU = 5 sec

Time for a single kernel to run on the CPU = 8 sec

Time for both, running concurrently = around 8 sec (but I'm seeing 13 sec)

I guess my question is about whether or not clEnqueueNDRangeKernel() will forward a kernel on to a second processor in a two-processor system if the first processor is already running a kernel.  Thanks!

Edit: A related question would be: is it possible to have two queues, or is there only one queue because there is only one global clEnqueueNDRangeKernel() ?

another thing. in release notes is stated that every element of structure must be aligned to float4. but i write struct and i get sizeof() that structure is 324 byte. which is exactly the sum size of data types in the struct. i tested it only on CPU. is on GPU some restriction in alignment?

second thing in release note is that array in struct is not supported. but for me it worked.

after searching my code for several months for a rather serious memory leak, I tried to reduce the problem to its core. It seems to happen either when a CommandQueue or a Kernel is being created, its arguments set, or during/after execution. I have attached a minimalistic C++ program, demonstrating this behaviour, to this message.

What it does is the following: during setup it creates a context, gets device handles, compiles a very simple increment kernel, creates a buffer of size 8 MByte and fills it with 0.

Then it does the following 100 times: get the command queue, create kernel from program, set buffer as kernel arg, enqueue kernel, wait for queue to finish. Allocation and deallocation is handled by the stack.

Afterwards the buffer, program, devices and the context are manually deallocated.

I made it so it halts

1. before setup, 
2. before allocating, executing, deallocing the queue and kernel 100 times, 
3. after that, and 
4. after I manually deallocate buffer, program, context etc.

if you look at task manager memory usage for the process, at the second and third halting points it should roughly be equal, and also at the last halting point it should be equal to the first. But it's not. Not at all, indeed! Here are my read-outs from windows task manager, when run on DEVICE_CPU:

1. 2,344 K
2. 22,708 K
3. 39,596 K
4. 31,380 K

so for 100 iterations, there were 3. - 2. = 16,888 K leaked. When I increase iteration count to 200, mem usage after kernel execution is 56,596 K, indicating a leak of 33,888 K!

300 iterations: 50,948 K leaked

400 iterations: 67,644 K leaked

This indicates a leak of ~169 K per Iteration.

For iteration counts over ~500, it fails during CommandQueue(), citing error code -6 -- Out of host memory.

When I halve the buffer size, the numbers don't change.

On DEVICE_GPU it leaks ~50 K per Iteration.

But maybe the problem is BKAC*, so please help me identify if there is something totally wrong with my memory allocation / deallocation pattern. Should I allocate queue and kernels only once during setup? I tried this in my production code once, but as soon as I had created the commandqueue handle the program refused to respond to input via the gui.

OS: Win7 x64

RAM: 4 GByte

Compiler: VC++ 2008

Devices: Athlon x64 CPU (1 GB reported), Juniper GPU (5770, 128 MB reported)

*) between keyboard and chair, i.e. me

I am new to using ATIStream SDK kit for OpenCL.

I am facing problems in installing the ATIStream SDK kit.

I have downloaded the 'atistream-1.4.0_beta-lnx32.tar.gzip'.

While I run the script 'atistream-brook-1.4.0_beta.i386.run'

I get the following error message:

-------------------------------------------------------------------

Installing package via RPM...

error: Failed dependencies:

libstdc++.so.6(GLIBCXX_3.4.5) is needed by atistream-brook-1.4.0_beta-1.i386

----------------------------------------------------------------------

On checking the system I find that libstdc++.so.6 exists in the system.

The system that I am using is a 32 bit machine and is running RedHat linux 2.6.9-34.EL.

Could anyone provide any help on this?

thanks

Mona

/tmp/OCL5IbHhj.cl(10): warning: null (zero) character in input line ignored

Any help would be much appreciated.

  i've got a hd radeon 5870 display card and latest 9.10 catalyst. Is there something i missed??

I has an ATI HD4870 RV770 and an intel Q8200 processor

I am on linux ubuntu 9.04 and Ati driver beta installed

Searching for CL_DEVICE_EXTENSIONS i get for the Q8200

"cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store"

and the RV770 return an empety string

Q8200 and RV770 has double support, but the extension "cl_khr_fp64" is not defined

I am doing something wrong ?

Can I execute OpenCL applications against an AMD GPU when I am accessing a Windows machine using Remote Desktop?

Thanks,

--Joe

$ ./AESEncryptDecrypt -d -i output.bmp -o unencryptedOutput.bmp

I get another 1-bit random noise .bmp back.  What am I missing?  Thanks!

-Chris

Can anyone in the know hint at a timeframe for GPU double precision via OpenCL?

Thanks,
Mike

http://www.browndeertechnology.com/docs/BDT_OpenCL_Tutorial_NBody.html

It uses libstdcl to simplify the host code, which was just mentioned in prev post.

Comments are welcome and appreciated.  Hope its useful.

-dar

Its free software and distributed under the LGPLv3 license.  Currently targets Linux.

More info can be found here:

http://www.browndeertechnology.com/stdcl.html

-dar

__kernel kern(__global int* a_arg){},

with argument a_arg initialized by

kern_kernel.setArg(0, a_buf),

where

cl::Buffer a_buf(...).

Is it possible to extract value of a_arg from a_buf?

I have built sample programs and try to run them on a Linux box with ATI Firestream 9250 card.

The fgrlx driver is installed on that machine.

The programs run, but they do not detect GPU.

Sample output:

~/ATI/ati-stream-sdk-v2.0-beta4-lnx64/samples/opencl/bin/x86_64$ (export LD_LIBRARY_PATH=../../../../lib/x86_64/:/opt/acml4.3.0-int64/gfortran64_int64/lib:/opt/acml4.3.0-int64/gfortran64_int64/lib:$LD_LIBRARY_PATH; ./BinarySearch)

Sorted Input

0 1 1 1 1 1 1 2 3 4 5 6 6 6 7 7 7 8 8 9 9 9 9 10 11 11 12 12 13 13 14 14 14 14 14 14 15 16 17 17 18 18 19 19 20 20 21 22 22 23 24 25 25 25 25 26 26 26 27 28 28 28 28 28 

For test only: Expires on Sun Feb 28 00:00:00 2010

No protocol specified

Unsupported GPU device; falling back to CPU ...

l = 10, u = 11, isfound = 1, fm = 5

How do I tell the program to use GPU?

Help would be greatly appreciated.

Zeljko

got some of the OpenCL examples running with my onboard 3200. Now tried an 4870 from XFX but  the device is not known.

Any hints? Is there a BIOS Problem with XFX Branding?

Greetings,

S.Matovic

Is my hardware broken, or what could be the problem??

1. Do you plan to release new extesions used by opencl ( and when ) ? Also if possible could you write short description what they do ( some are obvious ) ?

Here is the list:

calExtGetProc: extid=8007 name=calConfig
calExtGetProc: extid=8005 name=calCtxCreatePrivateCounter
calExtGetProc: extid=8005 name=calCtxConfigPrivateCounter
calExtGetProc: extid=8005 name=calCtxGetPrivateCounter
calExtGetProc: extid=8008 name=calResAllocView
calExtGetProc: extid=8008 name=calResQueryInfo
calExtGetProc: extid=8008 name=calResMemCopy
calExtGetProc: extid=8009 name=calCtxWaitForEvents ( is it blocking ? )
calExtGetProc: extid=800B name=calMemCopyRaw

2.  Are the devs going to implement LDS optimization for 4xxx family. Specifically I'm thinking about detecting if kernel writes to LDS match pattern "LDS[const1*p + const2]=value" ( where const2<const1 ). This would allow to use native LDS. And if memory access doesn't follow this pattern use global memory ( as it's done now ).

Probably most of the kernels will use this access pattern anyway and it would give huge speed advantage ( and probably some could be converted by programmers if they knew about this optimization ).

3. Are the devs planning to implement use of texture units for memory access ( 4xxx family ). Again the problem is detecting by compiler if memory reads follow the pattern "value=some_const_pointer_parameter[width*y + x]" ( where x<width and width is some value which could be computed by kernel ( or const or parameter ) ). As it's const it cannot be written and memory overlapping with other parameters could be detected at run time ( then we use current compiler code ).

This optimization is quite important for writing efficient kernels ( like matrix mul ).

4. CAL & 4xxx question. Access from/to memory by g[] variable ( global buffer ) generates code with UNCACHED flag. Is it possible to change it to CACHED ?

Example from some code:

07 TEX: ADDR(64) CNT(1)
              9  RD_SCATTER R0, DWORD_PTR[0+R0.x], ELEM_SIZE(3) UNCACHED

Documentation to 7xx ISA suggests CACHED flag should be available.

And one more thingy . If you can't answer some or all of this questions please write so .

Also I can add that for me opencl is unusable without points 2 & 3 ( i'm forced to use CAL or switch to other brand which could be less hassle ).

Hazeman

1) Detect the OpenCL.dll. Delay Load DLL is not enough. We need a method to detect if the OpenCL.dll is present on the system because, if not, any cl:evice or cl:latform will cause an ugly "Sorry, this program cannot find the module OpenCL.dll".

For windows, detect the Windows\system32\OpenCL.dll presency...

For linux/macos/solaris, I don't know

2) We need a way to see if the GPU is actually the primary display... so we can skip that card due to the 5 seconds watchdog or whatever. Just add a new CL_GPU_IS_PRIMARY_DISPLAY to the cl:evice.

thx

I've installed (and reinstalled) the Stream beta4 SDK, got the latest drivers for my card (v9.10), all the path variables are pointing to the correct places and the code seem to compile just fine.

I'm using Windows 7 64bit. I've tried to search the forums for anyone having a similar problem, but to no result.

Is it possible to do openCL computations on a RV610 graphics card ?

It's an on-board ATI radeon 3200.

With a test-program, I can compile just fine for the CPU, but any attempt to build the program for the GPU results in a "link failed". I tried some of the example, and also a tiny one:

__kernel void test(__global float yOut)
{
    //yOut = get_global_id(0);
}

Is the card just too cheap, or I am doing something wrong ?

I've attached the openCL info of the graphics-card. It does not have any extensions at all, but I'm not sure how important that is.

For example, in the attached code, the program will compile successfully if I remove the normalize call.  Furthermore, if I write my own normal function it will still fail.  I have tracked it down to the index variable(s) being involved in any way with a divide.  This only happens when it is compiled for the GPU by the way.

First, I think this must be a compiler bug.  Second, does anyone know how to normalize a vector without using a division?  I tried multiplying the vector by the length calculated using rsqrt to no avail.

My question is: is there a way to install the beta4 on a Windows machine that is lacking a GPU so that the CPU can be used as the ocl target device?

I'd like to use the beta4 runtime to target the CPU on a machine without a Radeon GPU, but the catalyst drivers won't install in the absence of a graphics card.

I've manually installed the aticalcl.dll, &c from a different machine, along with the demo applications.Now, when I try running one, I get:

$ ./RadixSort --device cpu
Error: clCreateContextFromType failed. Error code : CL_DEVICE_NOT_FOUND
For test only: Expires on Sun Feb 28 00:00:00 2010

MJH

For test only: Expires on Sun Feb 28 00:00:00 2010

CL Error: clBuildProgram (-11)

Build Log:

Link failed

Press any key to continue . . .

For test only: Expires on Sun Feb 28 00:00:00 2010

C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0x2c): undefined reference to `__vla_alloc'

C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0xba): undefined reference to `__vla_dealloc'

C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0x18e): undefined reference to `__vla_alloc'

C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0x21d): undefined reference to `__vla_dealloc'

CL Error: clBuildProgram (-11)

Build Log:

Compilation failed

Press any key to continue . . .

Sample code

global volatile flot4* v;

a1 = v[0]; <- this read is generated

a2 = v[0]; <- optimized out

a3 = v[0]; <- optimized out

The standard says "The type qualifiers const, restrict, volatile as defined by the C99 specification are supported". So I think this needs to be corrected.

I wanted to take the next step, and get code running on my Radeon Mobility 4830. I extracted ati-opencl-beta-driver-v2.0-beta4-lnx.zip. In the fglrx-8.67/ directory, I run "sudo ./ati-driver-installer-8.67-x86.x86_64.run". Everything appears to build and install successfully. The kernel module can load without trouble, but I get a black screen when X starts. It appears as though my computer has completely locked up.

If I remove the fglrx.ko kernel module, then X will start and not hang the system. I get the "For testing" AMD logo in the lower right corner of the screen. 2D performance is absolutely awful, and I get an error when I try to run my openCL applications. After the expiration message printout, I get an "X Error of failed request: ..." and then some information about failed request opcodes.

Am I missing something?

Nvidia samples AMD drivers

http://oscarbg.blogspot.com/2009/11/nvidia-sdk-with-amd-opencl.html

Nvidia samples Nvidia 195 driver

http://oscarbg.blogspot.com/2009/11/amd-opencl-samples-on-nvidia-195-opencl_05.html

AMD samples Nvidia 195 driver

http://oscarbg.blogspot.com/2009/11/amd-opencl-samples-on-nvidia-195-opencl.html

AMD and Nvidia SDK in VS 2010

http://oscarbg.blogspot.com/2009/11/optix-and-amd-opencl-sdk-visual-studio.html