OpenCL Kernel for Scalar Product with Atomic Operations

The final sum of the dotproduct example is implemented on CPU. This is a solution of Scalar Product (DotProduct) without final reduction on the host side. This example uses atomic operations.

/*
* sDOT OpenCL Kernel Function for Level 1 BLAS Dot Product dot<-xy * Author; Wendell Rodrigues 
* INRIA-Lille :: DaRT Team
*/
__kernel void sDOT(
__global const unsigned int N,
__global const float* X,
__global const float* Y,
__global float* DOT,
__global int* FLAG,
__local float* sdata
)
{
// get index into global data array
unsigned int tid = get_local_id(0);
unsigned int i = get_global_id(0);

sdata[tid] = (i<N) ? X[i]*Y[i] : 0;

if (i==0) {
DOT[0]=0;
*FLAG=0;
}

barrier(CLK_LOCAL_MEM_FENCE);


// do reduction in shared mem
for(unsigned int s=1; s < get_local_size(0); s *= 2)
{
  int index = 2 * s * tid;

  if (index < get_local_size(0))
  {
   sdata[index] += sdata[index + s];
  }

  barrier(CLK_LOCAL_MEM_FENCE);
}

// write result for this block to global mem
if (tid == 0) {
 while (atom_cmpxchg(FLAG,0,1)==1);
 DOT[0] += sdata[0];
 atom_cmpxchg(FLAG,1,0);
}

}

Modeling Challenges on OpenCL Code Generation

This week, I presented an overview of the integration of OpenCL and Gaspard2. In order to overcome the many challenges of model conception and transformations, we need study a new MoC (other than Array-OL), coalescent memory allocation and task distribute. The slides have two good examples of OpenCL applications and some questions about model conception and code generation.

PDF

OpenWF - The Standard for building composited windowing systems

From Khronos Group:

Embedded devices are increasingly expected to offer sophisticated user interfaces that combine rich graphics with multimedia content. Graphics and display hardware technologies have evolved to achieve these visuals with significantly higher efficiency than traditional CPUs, delivering greater performance, decreasing memory bandwidth usage and increasing battery life. Making use of this variety of hardware introduces fragmentation as software needs to be adapted to each hardware configuration.

A platform’s Hardware Abstraction Layer (HAL) for display and graphics technology allows the applications and middleware layers above to be deployed across a range of hardware without costly porting activities. OpenGL is an example of a graphics HAL that allows portable software to take advantage of a wide range of 3D hardware accelerators.

Windowing systems allow screens to be shared by multiple applications, ensuring that the graphics provided for each application’s window is sensibly merged onto the screen. This requires the graphics and display drivers to respect the intentions of the windowing system, which commonly means considerable OS-specific porting work on the part of the device manufacturer when moving to new hardware.

The OpenWF APIs provide an OS-independent and hardware-neutral foundation for building compositing systems, particularly suited to implementing windowing systems. OpenWF acts as a HAL to achieve composition of content and configuration of display devices. The interfaces are designed for use by a single user which could be a central windowing system or, in an application-specific system, may be the application itself.

http://www.khronos.org/openwf/

Conjugate Gradient and OpenCL

I've just finished a conjugate gradient implementation for OpenCL. It has not performance yet, but I'm working on this to fix the bugs and/or optimize the code.

Here you are a PDF that makes an overview on the subject.

Nvidia's Next Generation: Fermi - key architectural highlights

Third Generation Streaming Multiprocessor (SM)

• 32 CUDA cores per SM, 4x over GT200
• 8x the peak double precision floating point performance over GT200
• Dual Warp Scheduler simultaneously schedules and dispatches instructions from two independent warps
• 64 KB of RAM with a configurable partitioning of shared memory and L1 cache

Second Generation Parallel Thread Execution ISA

• Unified Address Space with Full C++ Support
• Optimized for OpenCL and DirectCompute
• Full IEEE 754-2008 32-bit and 64-bit precision
• Full 32-bit integer path with 64-bit extensions
• Memory access instructions to support transition to 64-bit addressing
• Improved Performance through Predication

Improved Memory Subsystem

• NVIDIA Parallel DataCache™ hierarchy with Configurable L1 and Unified L2
• Caches
• First GPU with ECC memory support
• Greatly improved atomic memory operation performance

NVIDIA GigaThread™ Engine

• 10x faster application context switching
• Concurrent kernel execution
• Out of Order thread block execution
• Dual overlapped memory transfer engines

more information: http://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf

ATI Stream Software Development Kit (SDK) v2.0 Beta Program

What’s New in v2.0-beta4

• First beta release of ATI Stream SDK with OpenCL™ GPU support.
• ATI Stream SDK v2.0 OpenCL™ is certified OpenCL™ 1.0 conformant by Khronos.
• Added Microsoft® Windows® 7 support.
• Added native Microsoft® Windows® 64-bit support.
• Float comparisons in kernels no longer produce a runtime error.
• Various other issues from previous v2.0 beta releases have been resolved.

More information: http://developer.amd.com/GPU/ATISTREAMSDKBETAPROGRAM/Pages/default.aspx

OpenCL BLAS - Makefile for MAC

Thanks to Mario Rometsch for a version of OpenCL BLAS Makefile for MacOS. You can download it on the SourceForge.

OpenCL BLAS Makefile for MacOS