70 views (last 30 days)
Michael on 30 Mar 2012
Commented: Matteo Piccio on 29 Jan 2021
Hi Forum,
I'm very new to GPU programming, and I'm trying to write a simple CUDA program to speed up calculations of the digamma function for large 2D matrices. I've adapted a C implementation of the digamma function (from the "lightspeed" Matlab toolbox) to run on the GPU. I then compile the .cu code to a .ptx file and instantiate the kernel in Matlab by calling
k = parallel.gpu.CUDAKernel('digamma_test.ptx','');
and then use
o = feval(k,matrix);
The function works great, providing huge speed-ups on large matrices. The problem is that after a bunch of feval calls I get an error when I try to instantiate the kernel or call feval: "CUDA_ERROR_LAUNCH_FAILED", with no more information than that. If I restart Matlab, I can get it to work again for about the same period of time.
My thoughts are that somehow I'm not managing my memory very well. I've never worked in a language where memory has to be explicitly managed like CUDA or C, so I'm sure I'm leaving something out. The problem is, I just don't know what.
Here's my kernel:
#include <math.h>
#include <stdlib.h>
#include <float.h>
/* this is directly from Minka's lightspeed toolbox*/
__global__ void digamma(double *y)
int idx_x = blockIdx.x;
int idx_y = blockIdx.y;
int offset = idx_x + idx_y * gridDim.x;
double x = y[offset];
double neginf = -INFINITY;
const double c = 12,
digamma1 = -0.57721566490153286,
trigamma1 = 1.6449340668482264365, /* pi^2/6 */
s = 1e-6,
s3 = 1./12,
s4 = 1./120,
s5 = 1./252,
s6 = 1./240,
s7 = 1./132,
s8 = 691./32760,
s9 = 1./12,
s10 = 3617./8160;
double result;
/* Illegal arguments */
if((x == neginf) || isnan(x)) {
y[offset] = NAN;
/* Singularities */
if((x <= 0) && (floor(x) == x)) {
y[offset] = neginf;
/* Negative values */
/* Use the reflection formula (Jeffrey 11.1.6):
* digamma(-x) = digamma(x+1) + pi*cot(pi*x)
* This is related to the identity
* digamma(-x) = digamma(x+1) - digamma(z) + digamma(1-z)
* where z is the fractional part of x
* For example:
* digamma(-3.1) = 1/3.1 + 1/2.1 + 1/1.1 + 1/0.1 + digamma(1-0.1)
* = digamma(4.1) - digamma(0.1) + digamma(1-0.1)
* Then we use
* digamma(1-z) - digamma(z) = pi*cot(pi*z)
if(x < 0) {
*p = digamma(p,1-x) + M_PI/tan(-M_PI*x);
/* Use Taylor series if argument <= S */
if(x <= s) {
y[offset] = digamma1 - 1/x + trigamma1*x;
/* Reduce to digamma(X + N) where (X + N) >= C */
result = 0;
while(x < c) {
result -= 1/x;
/* Use de Moivre's expansion if argument >= C */
/* This expansion can be computed in Maple via asympt(Psi(x),x) */
if(x >= c) {
double r = 1/x, t;
result += log(x) - 0.5*r;
r *= r;
#if 0
result -= r * (s3 - r * (s4 - r * (s5 - r * (s6 - r * s7))));
/* this version for lame compilers */
t = (s5 - r * (s6 - r * s7));
result -= r * (s3 - r * (s4 - r * t));
/* assign the result to the pointer*/
y[offset] = result;
  1 Comment
Joss Knight
Joss Knight on 17 Jan 2014
How are you setting your launch configuration? (ThreadBlockSize, GridSize). You are accessing your array in an unusual way using the blockIdx instead of threadIdx. If your array is just a vector it would be more normal for you to have:
offset = threadIdx.x + blockDim.x*blockIdx.x;
You're also going to need some sort of guard to prevent threads from accessing beyond the end of the array. The normal way to do this would be to be passing in the length of your array, or have it hard-coded into your kernel.
Note that if your kernel fails to launch it will put the device into an error state and no further kernels can be launched until the state is cleared. This can be done in MATLAB by resetting the device:
or in C code using

Sign in to comment.

Answers (2)

Edward on 14 Jan 2014
I am having a similar problem as well. My kernel "runs," but then all GPU functionality is lost. If I try to do anything involving the GPU, I am met with:
An unexpected error occurred during CUDA execution. The CUDA error was:
An unexpected error occurred during CUDA execution. The CUDA error was:
unspecified launch failure
Anyone have any ideas? Thanks!
Matteo Piccio
Matteo Piccio on 29 Jan 2021
@Marius Bledea I changed the settings as suggested but still having same issues. Do you have any other suggestion?
I'm using an NVIDIA GTX 770.
Thanks in advance

Sign in to comment.

fausto on 28 Jun 2013
Edited: fausto on 28 Jun 2013
i have a similar problem. my matlab won't even work for a short period of time. it's giving this error immediately.
did you solve it? does somebody have the answer to this?
  1 Comment
Tim DeWolf
Tim DeWolf on 11 Aug 2017
For me, I got this when a kernel hit the execution time limit on Windows. Solution for me is to break up the problem more so the kernel completes faster. I had to completely restart MATLAB as reset(gpuDevice) didn't work

Sign in to comment.


Community Treasure Hunt

Find the treasures in MATLAB Central and discover how the community can help you!

Start Hunting!