Function call_kernel_triple_update

src/library/blas/functor/gpu_dtrsm.cc:127–175

cl_int call_kernel_triple_update(const char* kernel_name, 
                                 const cl_program prg,
                                 const cl_command_queue queue,
                                 cl_mem A,
                                 unsigned int offA,
                                 cl_mem d_dinvA,
                                 int i,
                                 unsigned int lda,
                                 int M, 
                                 cl_event *event)
{
  cl_int err = 0;

  unsigned int m = M; 

  int npages = M/(i*2)+(M%(i*2)!=0);
  size_t globalLocal  [2] = { (i <= 32)?(i/4):16, 4};
  size_t globalThreads[2] = { (i/(globalLocal[0]*globalLocal[1]))* globalLocal[0],
				npages*(i/16) * globalLocal[1]};

  cl_kernel kernel = clCreateKernel(prg, kernel_name, &err);
  if (err != CL_SUCCESS) {
    //printf( "create kernel %s failed with %d\n", kernel_name, err );
    return err;
  }

  clSetKernelArg(kernel, 0, sizeof(cl_mem), &A);
  clSetKernelArg(kernel, 1, sizeof(unsigned int), &offA);
  clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_dinvA);
  clSetKernelArg(kernel, 3, sizeof(int), &i);
  clSetKernelArg(kernel, 4, sizeof(unsigned int), &lda);
  clSetKernelArg(kernel, 5, sizeof(int), &npages);
  clSetKernelArg(kernel, 6, sizeof(unsigned int), &m);

  err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL,
			       globalThreads, globalLocal , 
			       0, NULL, event);
  
 
  if (err != CL_SUCCESS) {
    clReleaseKernel(kernel);
    //printf( "execution of kernel %s failed with %d\n", kernel_name, err );
    return err;
  }

  err = clReleaseKernel(kernel);
  return err;

}