Khronos.org Message Boards

Confusion on output

Thu Feb 09, 2012 11:52 am

by mathuin (Posted Thu Feb 09, 2012 11:52 am)
I was able to solve my own problem. My main trouble was argument agreement -- don't send an integer when a float is expected, et cetera. The next step is to figure out how to handle arbitrarily large arrays, but that will be a different thread if necessary.

Here is the resulting code:

First the OpenCL code:

Code:

struct point {
  int x;
  int z;
};

__kernel void nearest(__global struct point *coords, __global int *values, __global struct point *base, __global int *output, const unsigned int lencoords) {
  // base index
  int g_dataset_id = get_global_id(0);

  int dmin = -1;
  int d, dx, dz;

  for (int i=0; i<lencoords; i++) {
    dx = coords[i].x-base[g_dataset_id].x;
    dx = dx * dx;
    dz = coords[i].z-base[g_dataset_id].z;
    dz = dz * dz;
    d = dx + dz;
    if (dmin == -1 || d < dmin) {
      output[g_dataset_id] = values[i];
      dmin = d;
    }
  }
}

Now the Python code:

Code:

import pyopencl as cl
import numpy
from itertools import product
from random import randint, uniform

class CL:
    def __init__(self):
        self.ctx = cl.create_some_context()
        self.queue = cl.CommandQueue(self.ctx)

    def loadProgram(self, filename):
        #read in the OpenCL source file as a string
        f = open(filename, 'r')
        fstr = "".join(f.readlines())
        #print fstr
        #create the program
        self.program = cl.Program(self.ctx, fstr).build()

    def popCorn(self):
        mf = cl.mem_flags

        #initialize client side (CPU) arrays
        self.xsize = 256
        self.zsize = 256
        self.lenbase = self.xsize * self.zsize
        self.lencoords = max(1, int(self.lenbase*0.05))
        self.coords = numpy.array([(randint(0, self.xsize-1),randint(0, self.xsize-1)) for elem in xrange(self.lencoords)], dtype=numpy.int32)
        self.values = numpy.array([uniform(1,5) for elem in xrange(self.lencoords)], dtype=numpy.int32)
        self.base = numpy.array([(index / self.zsize, index % self.zsize) for index in xrange(self.lenbase)], dtype=numpy.int32)
        self.output = numpy.zeros((self.lenbase), dtype=numpy.int32)

        #create OpenCL buffers
        self.coords_buf = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=self.coords)
        self.values_buf = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=self.values)
        self.base_buf = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=self.base)
        self.output_buf = cl.Buffer(self.ctx, mf.WRITE_ONLY, size=self.output.nbytes)

    def execute(self):
        self.program.nearest(self.queue, self.base.shape, None, self.coords_buf, self.values_buf, self.base_buf, self.output_buf, numpy.int32(self.lencoords))
        self.results = numpy.empty_like(self.output)
        cl.enqueue_read_buffer(self.queue, self.output_buf, self.results).wait()
        self.results.resize((self.zsize, self.xsize))

    def dumpit(self):
        for xind in xrange(self.xsize):
            line = ''
            for zind in xrange(self.zsize):
                line += str(int(self.results[xind,zind]))
            print line

            

if __name__ == "__main__":
    example = CL()
    example.loadProgram("nearest.cl")
    example.popCorn()
    example.execute()
    example.dumpit()

Cross-device bandwidth for discrete GPU (HD 5870)

Thu Feb 09, 2012 11:41 am

by Cadorino (Posted Thu Feb 09, 2012 11:41 am)
Hi,
I'm testing a system equipped with a Fusion A8-3850 and an HD 5870 gpu. I was planning to test the memory access bandwidth in the following cases:

1) The discrete GPU (HD 5870) reads from a buffer allocated in the host memory (CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY)
2) The integrated GPU (6550D) reads from a buffer allocated in the host memory (CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY)

Reads are performed linearly (each thread reads a fixed-size memory range starting from its own global index).

I was assuming that the result of the first test (discrete gpu) would have never been higher than the PCI-express bandwidth (approx 8GB/s), but I'm getting a bandwidth that is around 40 GB/s.
I'm checking the bandwidth by using both the GlobalMemoryTest sample shipped with the AMD SDK and a program written by myself. The results are very similar.

Can you explain me if it is (and why it is) possible to get a cross-domain (gpu->cpu) read bandwidth higher than the PCI one from a discrete GPU?.

Thank you very much!

CPU + GPU processing of same image

Thu Feb 09, 2012 11:01 am

by Photovore (Posted Thu Feb 09, 2012 11:01 am)
When I tried it, it was pretty easy and well worth the effort.

If it's working on CL_DEVICE_TYPE_GPU, the same code should work on the cpu if you specify CL_DEVICE_TYPE_CPU. Then make two contexts, two command queues, two programs, etc; just do everything twice, once for each. I read the data back to a big array in ram, so I just figure out the offset into that array when I read back the bottom half. I also send in a parameter so each kernel knows where it is. (I also do a 3-way, cpu+big gpu+little gpu, to try to get just a little more performance.)

What I don't know about is how you could merge the CPU half of the data up onto the graphics card itself, to display with opengl, without having to copy the gpu's data off and back on.

Real-time raytracing -- Is MIMD req'd?

Thu Feb 09, 2012 10:23 am

by Nick Wiggill (Posted Thu Feb 09, 2012 10:23 am)
(Rest the latter's soul.) Thanks for the input. These are new concepts to me (I've heard vaguely of branchless logic though I don't grasp the underlying mechanism yet. Any pointers on that would help.

This is why I mentioned books; I feel there are a lot of blanks I need to fill in.

Command queue goes invalid after kernel execution

Thu Feb 09, 2012 9:53 am

by Photovore (Posted Thu Feb 09, 2012 9:53 am)

Quote:

...invalid command queue usually means your code has crashed on the device. / Without the complete source, it's anyone's guess as to why.

Even _with_ the complete source, for my ~2000 lines of code I ultimately used the "méthode tedieuse" -- comment out huge chunks, verify that it runs. Uncomment out something, try again, till "command queue invalid" occurs. Of that chunk you've just uncommented, subdivide it and comment out bits of that. Find culprit and fix. On to next section. Tedious as I have many helper subroutines, and the strangest things seemed to set it off.
Of course, this was my first attempt at opencl; I had a project I'd worked on for decades that was a good candidate and I just dove into converting it. Made it work first, optimized later.

Quote:

(op on advice re: possibly running off the end of an array)

Make sure that it doesn't! One of those was the final gotcha for me; kernel would run fine for a random number of seconds, then command queue invalid. Was running off the end of an array only occasionally.

Is OpenCL can be used here?

Thu Feb 09, 2012 6:51 am

by notzed (Posted Thu Feb 09, 2012 6:51 am)

lloydkl wrote:

Thanks. The errors are only due to the lack of my knowledge in OpenCL. I hope I can learn and fix them.

My suspicion was also about the efficiency - as in my present code there is large amount of memory reallocation. As I mentioned earlier there is a large file, I have 'memory mapped" this file in the host (CPU) program (only file reading is needed). To pass this data (memory buffer) to OpenCL I again have created memory buffers for OpenCL, then copied data...

I didn't feel it as "efficient". I feel there must a better mechanism to "share" data between host and GPU without reallocation, isn't it?

Thanks,
Lloyd

I'm not sure why you'd need to 'reallocate', unless it's part of the algorithm.

No matter what you do the data still has to get to the cpu or the gpu, mmap still needs to read the disk, but it has no idea how the application will use the data so has to guess (read-ahead). If you're just streaming data, then you know the precise access patten so can read-ahead yourself both accurately and trivially: it should be possible to beat (or at least equal) mmap since the latency is from the disk access and not from the memcpy's.

Since the CPU is just shunting data around and GPU is doing the work, it's not like you're saving the cpu cycles for processing either.

For streaming in opencl you'd just allocate a few buffers and use them cyclically ('multi-buffer'), loading the next one while the current one is being sent to the gpu and so on.

This multi-buffer approach is very efficient and can hide the i/o and bus latencies. Assuming the processing takes longer than the PCI transfers, the cpu should just be queuing work and spending most of it's time waiting around for the GPU to finish the currently-oldest buffer queued. And if it takes more time than the disk i/o, the disk reads should be finished by the time the gpu is ready for the data too.

Obviously you have to 'copy' the data to the gpu device, since it uses different memory (unless you're using an APU, in which case the 'copy' functions do nothing).

Face/Object detection - are they OpenMAX IL components?

Thu Feb 09, 2012 6:16 am

by ngnir (Posted Thu Feb 09, 2012 6:16 am)
We are developing a face detection lib.

Is this kind of lib has a place in the OpenMAX IL?
If so, under which component class?
Which roles?

Thanks

Nir

doubts with work items and groups

Thu Feb 09, 2012 5:42 am

by mustang (Posted Thu Feb 09, 2012 5:42 am)
Thank you again!!! And regarding execution, I have an nvidia card (it is a 9400/ION), is it correct that each work group is executed by only one CUDA core (in test deviceQuery it says that it has 2*8=16 CUDA cores) but each CUDA core may execute more than 1 work group as there is no limits for the amount of them?
thank you very much for paying attention to my questions which I guess are very basic!!

Pablo

problem environment mapping using samplerCube

Thu Feb 09, 2012 2:31 am

by xcode_dave (Posted Thu Feb 09, 2012 2:31 am)
Hi everyone,

I'm having a strange issue with environment mapping by using samplerCube for reflections (on an iPhone 4).

I don't think the issue is with my shader code because I've tested the output without the samplerCube (by using the reflection vector as the gl_FragColor), and it seems to output as I'd expect.

When I try using the samplerCube though, all fragments are just output as black colour.

Here is my shader code, first of all the vertex shader:

Code:

//  Shader.vsh

attribute vec4 position;
attribute vec3 normal;

uniform mat4 Projection;
uniform mat4 Modelview;
uniform mat3 Model;
uniform vec3 EyePosition;

varying vec3 ReflectDir;

void main()
{
    gl_Position = Projection * Modelview * position;
    // Compute eye direction in object space:
    highp vec3 eyeDir = normalize(position.xyz - EyePosition);

    // Reflect eye direction over normal and transform to world space:
    ReflectDir = Model * reflect(eyeDir, normal);
}

And here is my fragment shader:

Code:

//  Shader.fsh
varying highp vec3 ReflectDir;
uniform samplerCube cubeMap;

void main()
{
    gl_FragColor = vec4(ReflectDir.x, ReflectDir.y, ReflectDir.z, 1.0);
    //gl_FragColor = textureCube(cubeMap, ReflectDir);
}

In my fragment shader, you can see I've commented out the line causing the issue which tries to use the reflection direction vector to sample the cube map.
Above it, you can see my debug code which is using the reflection direction as a colour (to visually debug).

I get no errors during build or run, just the wrong output.

To load my cube map, I'm using GLKit for convenience. Here is my code for that:

Code:

//setup the cube map
        NSDictionary *options = [NSDictionary dictionaryWithObject:[NSNumber numberWithBool:YES] forKey:GLKTextureLoaderGenerateMipmaps];
        NSString *rightTex = [[NSBundle mainBundle] pathForResource:@"right" ofType:@"png"];
        NSString *leftTex = [[NSBundle mainBundle] pathForResource:@"left" ofType:@"png"];
        NSString *topTex = [[NSBundle mainBundle] pathForResource:@"top" ofType:@"png"];
        NSString *bottomTex = [[NSBundle mainBundle] pathForResource:@"down" ofType:@"png"];
        NSString *frontTex = [[NSBundle mainBundle] pathForResource:@"front" ofType:@"png"];
        NSString *backTex = [[NSBundle mainBundle] pathForResource:@"back" ofType:@"png"];
        NSError *error;
        cubeMapInfo = [GLKTextureLoader cubeMapWithContentsOfFiles:[NSArray arrayWithObjects:rightTex, leftTex, topTex, bottomTex, frontTex, backTex, nil] options:options error:&error];
        if (error!=nil){
            NSLog(@"ERROR: %@", error.localizedDescription);
            NSLog(@"Failure reason: %@", error.description);
            NSLog(@"Error code: %i", error.code);
        } else {
            NSLog(@"No error!");
            [cubeMapInfo retain];
            NSLog(@"Cube map name: %i", cubeMapInfo.name);
            glUniform1i(uniforms[UNIFORM_CUBE_MAP], cubeMapInfo.name);
        }

I'm not sure how I can debug further or what to do next. If anyone has any idea then your help would be much appreciated.

Thanks!

OpenWF - C on Android

Wed Feb 08, 2012 11:54 pm

by lremes (Posted Wed Feb 08, 2012 11:54 pm)
I'm not aware of any public attempts to port and integrate OpenWF to Android platform.
I haven't taken a closer look, but I assume you would somehow need to integrate OpenWF below SurfaceFlinger in the graphics stack.

Please note that the Sample Implementation does not use HW acceleration at all. If you want to HW accelerate the composition, for example, using OpenGL ES, you need to write a new backend.

WFD functions ported to Intel Sandy Bridge chip

Wed Feb 08, 2012 11:48 pm

by lremes (Posted Wed Feb 08, 2012 11:48 pm)
The public reference implementation is available at http://www.khronos.org/registry/wf/

I have not heard that anyone would yet have a public port of WFD for Sandy Bridge.
Any platform specific ports are usually property of the vendors.

The Sample Implementation contains a generic porting layer implementation for Linux that uses X11/SDL, but it’s just an example. As WFD is an abstraction of display controller and you probably want to get rid of any external dependencies, using that as the baseline does not work.

Can Native Kernels Enqueue Non-Native Kernels?

Wed Feb 08, 2012 5:58 pm

by sean.settle (Posted Wed Feb 08, 2012 5:58 pm)
I know non-native kernels cannot enqueue non-native kernels, but what about native kernels enqueueing non-native kernels? Say you have a quad-core CPU and four GPUs, then partition the CPU info four single-core devices, each of which would then control a distinct GPU.

When will out-of-order command queues be supported in almost all platforms?

kernel freezes video output

Wed Feb 08, 2012 4:52 pm

by Yours3lf (Posted Wed Feb 08, 2012 4:52 pm)
Hi,

I wrote a tile based deferred shading kernel, it compiles, but it makes the video output freeze after a few seconds of running with ~70FPS. Only the graphics card is hanging, the system is still responsive through ssh. I'm using Linux Mint 12.1 64 bit with Catalyst 12.1 64 bit, and APP SDK 2.6

here's the kernel:

Code:

    __constant float far = -10000.0f; //far plane distance  
    __constant float near = -1.0f; //near plane distance  
    __constant sampler_t the_sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;  
    __constant float cutoff = 0.25f; //0.005f  
    __constant int attenuation_type = 0; //linear or full attenuation?  
      
    float my_abs( float var ) //these floating point operations aren't supported in opencl 1.1  
    {  
      if ( var < 0 )  
      {  
        return -var;  
      }  
      else  
      {  
        return var;  
      }  
    }  
      
    float my_mix( float x, float y, float weigth )  
    {  
      return x * ( 1.0f - weigth ) + y * weigth;  
    }  
      
    float3 my_reflect( float3 incident, float3 normal )  
    {  
      return incident - 2.0f * dot( normal, incident ) * normal;  
    }  
      
    float3 decode_normals_spheremap( float4 n ) //decode normals from spheremap encoding  
    {  
      float4 nn = n * ( float4 )( 2.0f, 2.0f, 0.0f, 0.0f ) + ( float4 )( -1.0f, -1.0f, 1.0f, -1.0f );  
      float l = dot( nn.xyz, -nn.xyw );  
      nn.z = l;  
      nn.xy *= sqrt( l );  
      return nn.xyz * 2.0f + ( float3 )( 0.0f, 0.0f, -1.0f );  
    }  
      
    float3 decode_linear_depth( float4 linear_depth, float4 position ) //decode linear depth into view space position  
    {  
      return ( float3 )( position.xy * ( far / position.z ), far ) * linear_depth.x;  
    }  
      
    __kernel void main( __read_only image2d_t albedo, //diffuse surface color from the g-buffer  
                        __read_only image2d_t normals, //normals encoded using spheremap encoding  
                        __read_only image2d_t depth, //linear depth  
                        __write_only image2d_t result, //the output buffer that stores lighting data  
                        __global const float* far_plane, //the lower left and upper right corners of the far plane  
                        __global const float* in_view_pos, //view space camera position  
                        __global const float* in_lights, //1024 lights {light pos[3], diffuse_color[3], radius[1], specular intensity[1] }  
                        __global const float* in_num_of_lights, //number of incoming lights (1024)  
                        __global const float* in_projection_matrix ) //the projection matrix is used for frustum culling  
    {  
      /* 
       * Per pixel calculations (global) 
       */  
      
      int2 coords = ( int2 )( get_global_id( 0 ), get_global_id( 1 ) );  
      float4 raw_albedo = read_imagef( albedo, the_sampler, coords );  
      float4 raw_normal = read_imagef( normals, the_sampler, coords ); //this will store the decoded normals  
      
      float4 out_color = ( float4 )( 0.0f ); //this will store the resulting color  
      out_color.w = 1.0f;  
      
      float4 raw_depth; //this will store the decoded view space position  
      int2 global_size = ( int2 )( get_global_size( 0 ), get_global_size( 1 ) );  
      
      /* 
       * Per tile data 
       */  
      
      float3 view_pos; //this will store the view space position (uniform among the workgroups, but stored as local for speedup)  
      int num_of_lights; //num of lights (same here)  
      
      int2 local_coords;  
      int2 local_size;  
      int workgroup_index;  
      
      __local int tile_lights[1024]; //index of the lights visible per tile  
      __local int num_of_tile_lights; //number of lights per tile  
      
      float2 tile_scale; //used for calculating frustum culling, taken from Intel's sample  
      float2 tile_bias;  
      
      float4 column_1;  
      float4 column_2;  
      float4 column_4;  
      
      float4 frustum_planes[6];  
      
      /* 
       * Check for skybox 
       */  
      
      bool early_rejection = ( length( raw_normal.xy ) == 0.0f );  
      
      if ( early_rejection )  
      {  
        out_color = raw_albedo;  
      }  
      else  
      {  
        local_coords = ( int2 )( get_local_id( 0 ), get_local_id( 1 ) );  
        local_size = ( int2 )( get_local_size( 0 ), get_local_size( 1 ) );  
      
        workgroup_index = local_coords.y * local_size.x + local_coords.x;  
      
        float4 ll, ur;  
      
        raw_depth = read_imagef( depth, the_sampler, coords );  
      
        ll = ( float4 )( far_plane[0], far_plane[1], far_plane[2], 1.0f );  
        ur = ( float4 )( far_plane[3], far_plane[4], far_plane[5], 1.0f );  
      
        //texture coordinate [0...1] for input processing  
        float2 texel = ( float2 )(( float )( coords.x ) / ( float )( global_size.x ), ( float )( coords.y ) / ( float )( global_size.y ) );  
      
        raw_depth.xyz = decode_linear_depth( raw_depth, ( float4 )( my_mix( ll.x, ur.x, texel.x ), my_mix( ll.y, ur.y, texel.y ), ll.z, 1.0f ) );  
        raw_normal.xyz = decode_normals_spheremap( raw_normal );  
      
        view_pos = vload3( 0, in_view_pos );  
        num_of_lights = ( int )in_num_of_lights[0];  
      
        //I'm not sure if the maths here is correct due to OpenGL and DirectX using different matrices, but this shouldnt matter  
        tile_scale = ( float2 )( global_size.x, global_size.y ) * ( 1.0f / ( float )( 2.0f * local_size.x ) );  
        tile_bias = tile_scale - ( float2 )( local_coords.x, local_coords.y );  
      
        column_1 = ( float4 )( in_projection_matrix[5] * tile_scale.x, 0.0f, tile_bias.x, 0.0f );  
        column_2 = ( float4 )( 0.0f, -in_projection_matrix[10] * tile_scale.y, tile_bias.y, 0.0f );  
        column_4 = ( float4 )( 0.0f, 0.0f, 1.0f, 0.0f );  
      
        frustum_planes[0] = column_4 - column_1;  
        frustum_planes[1] = column_4 + column_1;  
        frustum_planes[2] = column_4 - column_2;  
        frustum_planes[3] = column_4 + column_2;  
        frustum_planes[4] = ( float4 )( 0.0f, 0.0f, -1.0f, near );  
        frustum_planes[5] = ( float4 )( 0.0f, 0.0f, 1.0f, far );  
      
        for ( int c = 0; c < 4; c++ ) //normalize frustum plane normals  
        {  
          frustum_planes[c].xyz *= 1.0f / length( frustum_planes[c].xyz );  
        }  
      
        /* 
        * Per workgroup (tile) calculations (local) 
        */  
      
        if ( workgroup_index == 0 )  
        {  
          num_of_tile_lights = 0;  
        }  
      }  
      
      barrier( CLK_LOCAL_MEM_FENCE );  
      
      if ( !early_rejection )  
      {  
        for ( int c = workgroup_index; c < num_of_lights; c += local_size.x * local_size.y ) //cull each light per tile, each thread in a tile processes one light  
        {  
          if ( c < num_of_lights )  
          {  
            bool in_frustum = true;  
            float attenuation_end = 0.0f;  
      
            if ( attenuation_type == 0 )  
            {  
              attenuation_end = ( float )( in_lights[c * 8 + 6] ) / ( float )( cutoff ); //radius / cutoff  
            }  
            else  
            {  
              attenuation_end = ( float )( in_lights[c * 8 + 6] ); //radius  
            }  
      
            for ( int d = 0; d < 6; d++ ) //cull each light based on the distance where it will shine and the frustum defined by the tile  
            {  
              float e = dot( frustum_planes[d], ( float4 )( in_lights[c * 8 + 0], in_lights[c * 8 + 1], in_lights[c * 8 + 2], 1.0f ) );  
              in_frustum = in_frustum && ( e >= -attenuation_end );  
            }  
      
            if ( in_frustum ) //if the light is in the frustum, then store its index (if I comment this out, the kernel runs, but doesn't cull lights)  
            {  
              int index = atomic_inc( &num_of_tile_lights );  
              tile_lights[index] = c;  
              if ( num_of_tile_lights != 0 )  
              {  
                out_color = ( float4 )( 0.0f, 1.0f, 0.0f, 1.0f );  
              }  
            }  
          }  
        }  
      }  
      
      barrier( CLK_LOCAL_MEM_FENCE );  
      
      /* 
       * Per light calculations 
       */  
      
      if ( !early_rejection )  
      {  
        for ( int c = 0; c < num_of_tile_lights; c++ ) //draw each light per tile  
        {  
          int index = tile_lights[c]; //get back the light index  
          float3 light_pos = ( float3 )( in_lights[index * 8 + 0], in_lights[index * 8 + 1], in_lights[index * 8 + 2] ); //gather light data using the index  
          float light_radius = in_lights[index * 8 + 6];  
      
          //calculate blinn-phong lighting with custom attenuation  
          float3 light_dir = light_pos - raw_depth.xyz;  
          float distance = length( light_dir );  
          light_dir /= distance;  
      
          float coeff, attenuation;  
      
          if ( attenuation_type == 0 )  
          {  
            coeff = max( distance - light_radius, 0.0f ) / light_radius + 1.0f;  
            attenuation = max(( 1.0f / ( coeff * coeff ) - cutoff ) / ( 1.0f - cutoff ), 0.0f );  
          }  
          else  
          {  
            attenuation = ( light_radius - distance ) / ( light_radius * 0.01f ) * 0.01f;  
          }  
      
          if ( attenuation > 0.0f )  
          {  
            float3 light_diffuse_color = ( float3 )( in_lights[index * 8 + 3], in_lights[index * 8 + 4], in_lights[index * 8 + 5] );  
            float light_specular_power = ( float )in_lights[index * 8 + 3];  
            float3 view_dir = normalize( view_pos - raw_depth.xyz );  
      
            float3 half_vector = ( light_dir + view_dir ) * 0.5f;  
            float n_dot_l = max( dot( raw_normal.xyz, light_dir ), 0.0f );  
      
            out_color.xyz += raw_albedo.xyz * light_diffuse_color * n_dot_l * attenuation;  
            float n_dot_h = pow( max( dot( raw_normal.xyz, half_vector ), 0.0f ), light_specular_power );  
            out_color.xyz += light_diffuse_color * n_dot_h * attenuation;  
          }   
        }  
      }  
      
      if ( coords.x < global_size.x && coords.y < global_size.y )  
      {  
        write_imagef( result, coords, out_color ); //write the calculated light data to the result buffer (texture)  
      }  
    }   

best regards,
Yours3!f

multiple OpenCL applications and global memory effects

Wed Feb 08, 2012 4:07 pm

by david.garcia (Posted Wed Feb 08, 2012 4:07 pm)
Different OpenCL contexts are isolated from each other in a similar manner that different processes are isolated.

Note I'm not saying that different OpenCL contexts are the same thing as different processes.

Why can kernels take __local pointer arguments?

Wed Feb 08, 2012 3:51 am

by homemade-jam (Posted Wed Feb 08, 2012 3:51 am)
Great thanks. I was clarifying since the docs I have read are very clear on what it can do but not on what it can't do, so just wanted to clarify.

Thanks.

Required atomic built-in functions

Tue Feb 07, 2012 4:16 pm

by notzed (Posted Tue Feb 07, 2012 4:16 pm)

sean.settle wrote:

I think I just found part of my answer. In OpenCL 1.0 basic atomic functions were optional, but as of OpenCL 1.1 they are required. If that is wrong please correct me.

Well, atomic functions pretty much ruins my hope of making an abstract device that integrates several similar devices. I could only think of how to do it without any atomic functions.

Any reasoning behind making it obligatory in OpenCL 1.1?

OpenCL is a software abstraction: you can implement atomics however you want, they just have to honour the contract.

e.g. for many devices you could break the kernels up into atomically bounded sections and run the kernel parts separately and then synchronise on the host.

I'm not saying it would be efficient, but I mean what do you really expect to be able to do anyway? The atomic operations require very specific specialised hardware in order to run fast, and without that you will have no choice but to resort to *host-based* software.

Global atomics are so slow on AMD hardware for example I wouldn't use them except for very rarely-executed code (i.e. it's possible calling the host already), so a high overhead is already expected. But they have global counters implemented in hardware to get around that ...

Re your earlier query there's nothing to say a research project implemented the full specification in the first place. It is possible to do without atomics entirely, at a cost of memory and extra processing steps.

Amalgamating different hardware with different performance characterstics will be a challenge! Often different hardware requires a different coding approach, it runs at a different speed, and so on: managing all that scheduling and keeping the memory close to the right kernels will be difficult.

Huge kernel overhead on Mac

Tue Feb 07, 2012 4:03 pm

by notzed (Posted Tue Feb 07, 2012 4:03 pm)

yoavhacohen wrote:

That makes sense, thanks.
I think that the specification should allow prevention of such lazy operations by adding flags to the kernel and the buffer constructors.
It might be the case that the application can do the allocation asyncroniuously but not the running of the kernel itself. In such cases, the lazy approach is a waste of time.

How is the lazy approach a waste of time? It only happens once, it still has to happen once no matter what happens. If you're doing any micro-benchmarks it is a given that you cannot get reliable results if you do not let the system `warm up' first - i.e. do a couple of dummy runs.

BTW the 'lazy allocation' is at the operating system level, and beyond such a specification's scope. Although a unix process can be given a big virtual address space, those pages do not exist until they are accessed. This is not something a driver could change.

picking

Tue Feb 07, 2012 4:03 pm

by gibubu (Posted Tue Feb 07, 2012 4:03 pm)
On user click of a particular object, I want to find the position of that click within the texture,

After that, I'll then update that texture with a red dot at that position.

I only need help with finding the position of that click.

Memory Allocation and time collection

Tue Feb 07, 2012 3:58 pm

by notzed (Posted Tue Feb 07, 2012 3:58 pm)

homemade-jam wrote:

notzed wrote:

Private memory is just the same as global memory if it isn't in a register. Its in a register if it can fit, and if it is an array with fixed indexing (or indexing that can be calculated at compile time).

Anyway this is pretty fundamental computer architecture: yes, if you can always use the fastest memory, the the next, and so on. So that means registers, then shared memory, then private/global memory.

Ok so I'm slightly confused now. Private memory is only accessible within a workitem; ie. you can't reach it from other workitems? Therefore, how does OpenCL know that it can be put into private memory? When you say shared memory, do you mean local memory? Or are we getting muddled with CUDA/OpenCL?

In my use case, caching isn't of much use since I only use each piece of data once.

Ahh yeah sorry, shared==local: apart from picking it up from nvidia's 'ported' opencl docs, it is the only memory shared amongst work items.

Private memory is private yes (oddly enough ...). If it fits/can be it goes into a register. The compiler knows it can go into private as private is the default qualifier for variable declarations. local and global are the other qualifiers.

Well you implied you were looping on the data in your original question.

But like i said, it depends on the problem. I'm sure you'll end up having more than a single problem to solve though, so you will eventually come across a use for it. Local memory can be used for more than a cache too, for example to re-arrange scatter-gather global memory requests for more efficient lookup, or particularly to communicate partial results amongst work-items (for reductions). But if you are only reading data once, and it's only being read sequentially (i.e. fully coalesced), then local memory just adds overhead.

OpenCL device problems

Tue Feb 07, 2012 3:29 pm

by wishgranter (Posted Tue Feb 07, 2012 3:29 pm)
have unistalled the AMD CPU driver and the app see my Nvidia GPU back...... Thanx

Example for Random Number Generator?

Tue Feb 07, 2012 10:04 am

by yoavhacohen (Posted Tue Feb 07, 2012 10:04 am)

Quote:

If so, I can't really make it fast without consuming more memory and have a states vector per work-item, right?

Not right: it's possible to save the memory overhead using a counter-based RNG:
http://www.openclblog.com/2011/11/gpus- ... ation.html

Drawing masked sprites (textured quad with alpha channel)

Tue Feb 07, 2012 9:39 am

by lucafr (Posted Tue Feb 07, 2012 9:39 am)
Hi, i would like to know how to draw a sprite (rendered using a textured quad) on the screen in such a way that some parts of it are not rendered at all (masked).

For example, suppose i have a wall with a window at its center. The window is open and i want to render a sprite inside the window. Suppose that the sprite is animated: the character is walking from left to right. I want to appear behind the window gradually.

The easy way is to use glScissor. Even if the window is not perfectly rectangular, i can use, for glScissor, a slightly bigger rectangle.

But what i i want to use a non rectangular mask (and even a non polygonal but complex mask)?

OpenMAX AL DTV sample implementation

Tue Feb 07, 2012 3:24 am

by francois (Posted Tue Feb 07, 2012 3:24 am)
Hi

The DTV extension looks very promising.
Is there a plan for a sample implementation ?

Thanks
Francois

Add camera functions

Tue Feb 07, 2012 12:32 am

by davepetalcurin (Posted Tue Feb 07, 2012 12:32 am)
Hi, maybe you can demonstrate it or further describe what you wanted so we can figure it out.

PTX file to work with clCreateProgramWithBinary

Mon Feb 06, 2012 3:44 pm

by luizdrumond (Posted Mon Feb 06, 2012 3:44 pm)
Hi,

My question is very simple.

I want to use clCreateProgramWithBinary, however i need the .ptx file to work with it.
Where is the file? I don't can find it.
I already used clCreateProgramWithSource before to try to use clCreateProgramWithBinary.

I work with VS2010.

In AMD SKD the file is also .ptx ?

Thanks,

Luiz.