cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

tball
Journeyman III

OpenCL 2.01, can't write to char pointers

opencl, fglrx, uchar, pointers

Hello

I have a problem witing to a char array, when using the GPU.

See my code attached.

I use the following:

Ubuntu 9.10

Fglrx 10.1/10.2

AMD stream SDK 2.01.

My output picture is somewhat distorted. See my attached picture. On a nvidia gfx card and using the cpu, there is no problem "copying" the uchar image to another.

I use the correct extension, to be able to write to an uchar pointer array.

#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable __kernel void hello(__global uchar * input, __global uchar * output, __global int * widthstep) { int tidx = get_global_id(0); int tidy = get_global_id(1); output[tidx + tidy*widthstep[0]] = input[tidx + tidy*widthstep[0]]; }

0 Likes
19 Replies
nou
Exemplar

cl_khr_byte_addressable_store is not supported at the moment. but you can try set enviroment variable GPU_BYTE_ADDRESSABLE_STORE=1

0 Likes
tball
Journeyman III

Thx for the quick response.

well running:

export = GPU_BYTE_ADDRESSABLE_STORE=1

./opencltest

Doesn't work. It does exactly the same.

0 Likes

no thats wrong.

export GPU_BYTE_ADDRESSABLE_STORE=1

0 Likes
tball
Journeyman III

Ups. Well that was what I did. 🙂

It doesn't work unfortunately.

0 Likes

it is undocumented feature. so it no big suprise it didnt work. but most likely it will work only on 5xxx series card. AMD for 4xxx series card do not plan any extension.

0 Likes
tball
Journeyman III

Thx for your help.

I will try convert my frame to floats instead.

0 Likes
tball
Journeyman III

Hmm I am trying to do a uchar to float typecast. See the code.

I get a lot of corruption in the output image again.

Should I just stay completely away from uchar and opencl?

#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable __kernel void hello(__global uchar * input, __global float * output, __global int * widthstep) { int tidx = get_global_id(0); int tidy = get_global_id(1); output[tidx + tidy*widthstep[0]] = (float)input[tidx + tidy*widthstep[0]]/255.0f; }

0 Likes

tball,
This should be fixed in our upcoming release, I get correct results with an internal build.
0 Likes

Thx Micah.

Are you going to support the cl_khr_byte_addressable_store part, or is it just the uchar to float typecast?

If cl_khr_byte_addressable_store is going to be supported, will the 4xxx series work with that extension?

0 Likes

tball,
The 4XXX series does not have the requisite hardware to support byte addressable store, so it will never be supported.
0 Likes

I tested writing into char pointer  and it worked in HD5870.

 

0 Likes

I can't write into char or short arrays. Neither private nor local memory.

 

 

System:

ATI 5770, AMD X2

Catalyst 10.5

SDK 2.1

Ubuntu 9.04 x86_64

0 Likes

smatovic,

Could you provide a testcase?

0 Likes

inside the kernel, integer works:

int board[82];

board[0] = 111;

 

but char not

char board[82];

board[0] = 111;

 

clBuildProgram fails.


0 Likes

Please check the buildlog to get the exact error message(you can also use SKA) and post it.

0 Likes

clBuildProgram returns -11.

Is there another way to get an error message?

I thought SKA is only available on Windows.

0 Likes

./my_kernels.cl(108): error: write to < 32 bits via pointer not allowed
          unless cl_khr_byte_addressable_store is enabled
      board[0] = 111;
      ^

0 Likes

ok, i think i got it:

 

#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable

 

char pointers are working.

0 Likes

clGetProgramBuildInfo can be used to return the verbose build log.

 

0 Likes