YA unspecified launch failure

cudapath · November 13, 2010, 2:44am

I’ve gathered that an “unspecified launch failure” is generic, but often caused by errant memory access. Below I have some partial code, and I’ve isolated the problem to one line:

extern __shared__ signed int sharedMem[];

// [... some code ..]

__global__ kernel(signed int *timeSeriesArray, [...]) {

  signed int * globalReferenceTimeSeries = (signed int *)(timeSeriesArray + referenceTimeSeriesIndex);;

  signed int * sharedReferenceTimeSeries = (signed int *)(sharedMem);

  // [... some code ..]

  for(unsigned int i = 0; i < timeSeriesLength; i++)

    // Problem line below:

    sharedReferenceTimeSeries[i] = globalReferenceTimeSeries[i];

}

But when I do

for(unsigned int i = 0; i < timeSeriesLength; i++)

  signed int foo = globalReferenceTimeSeries[i];

OR

for(unsigned int i = 0; i < timeSeriesLength; i++)

  sharedReferenceTimeSeries[i] = i;

OR even

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  signed int foo = globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = i;

OR

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  signed int foo = 12;

  sharedReferenceTimeSeries[i] = foo;

No problem. Then I do:

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  signed int foo = globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = foo;

And I’m back to the “unspecified launch failure” error with kernel termination.

Any ideas on:

How this is even possible (my C/C++ is 15 years old, so consider me a n00b if it’s an obvious pointer issue or some such thing)
How might I address this?

Thanks,

-CP

cudapath · November 13, 2010, 2:44am

I’ve gathered that an “unspecified launch failure” is generic, but often caused by errant memory access. Below I have some partial code, and I’ve isolated the problem to one line:

extern __shared__ signed int sharedMem[];

// [... some code ..]

__global__ kernel(signed int *timeSeriesArray, [...]) {

  signed int * globalReferenceTimeSeries = (signed int *)(timeSeriesArray + referenceTimeSeriesIndex);;

  signed int * sharedReferenceTimeSeries = (signed int *)(sharedMem);

  // [... some code ..]

  for(unsigned int i = 0; i < timeSeriesLength; i++)

    // Problem line below:

    sharedReferenceTimeSeries[i] = globalReferenceTimeSeries[i];

}

But when I do

for(unsigned int i = 0; i < timeSeriesLength; i++)

  signed int foo = globalReferenceTimeSeries[i];

OR

for(unsigned int i = 0; i < timeSeriesLength; i++)

  sharedReferenceTimeSeries[i] = i;

OR even

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  signed int foo = globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = i;

OR

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  signed int foo = 12;

  sharedReferenceTimeSeries[i] = foo;

No problem. Then I do:

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  signed int foo = globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = foo;

And I’m back to the “unspecified launch failure” error with kernel termination.

Any ideas on:

How this is even possible (my C/C++ is 15 years old, so consider me a n00b if it’s an obvious pointer issue or some such thing)
How might I address this?

Thanks,

-CP

tera · November 13, 2010, 3:31am

In all but the first and last case the compiler will optimize away the access to [font=“Courier New”]globalReferenceTimeSeries[i][/font] because its value is not actually used. So this strongly hints at a problem with [font=“Courier New”]timeSeriesArray[/font] or [font=“Courier New”]referenceTimeSeriesIndex[/font].

tera · November 13, 2010, 3:31am

In all but the first and last case the compiler will optimize away the access to [font=“Courier New”]globalReferenceTimeSeries[i][/font] because its value is not actually used. So this strongly hints at a problem with [font=“Courier New”]timeSeriesArray[/font] or [font=“Courier New”]referenceTimeSeriesIndex[/font].

cudapath · November 13, 2010, 3:52am

I tested

signed int foo;

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  foo = globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = i;

}

As well, which worked. Since “foo” is used later, I would expect the compiler not to be able to optimize away globalReferenceTimeSeries[i]. But

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  sharedReferenceTimeSeries[0] = globalReferenceTimeSeries[0];

}

Fails. So, there’s definitely something I’m missing here. I’m trying to reproduce the error in a “clean” kernel, but I am really at a loss here.

cudapath · November 13, 2010, 3:52am

I tested

signed int foo;

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  foo = globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = i;

}

As well, which worked. Since “foo” is used later, I would expect the compiler not to be able to optimize away globalReferenceTimeSeries[i]. But

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  sharedReferenceTimeSeries[0] = globalReferenceTimeSeries[0];

}

Fails. So, there’s definitely something I’m missing here. I’m trying to reproduce the error in a “clean” kernel, but I am really at a loss here.

cudapath · November 13, 2010, 4:06am

Thanks tera, you were right. Very right.

globalReferenceTimeSeries was a pointer to host memory in the kernel in question.
“foo” was used later, but not with the assigned the value of globalReferenceTimeSeries[i]. That’s some sophisticated optimization.

Is there a way to get more verbose warnings (e.g., using a host pointer on a device)?

cudapath · November 13, 2010, 4:06am

Thanks tera, you were right. Very right.

globalReferenceTimeSeries was a pointer to host memory in the kernel in question.
“foo” was used later, but not with the assigned the value of globalReferenceTimeSeries[i]. That’s some sophisticated optimization.

Is there a way to get more verbose warnings (e.g., using a host pointer on a device)?

tera · November 13, 2010, 4:15am

If [font=“Courier New”]foo[/font] is used later, then probably all but the last read of [font=“Courier New”]globalReferenceTimeSeries[i][/font] are optimized away. This would point to a (usually less common) problem at the beginning of the array.

You can verify that theory by checking that

signed int foo = 0;

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  foo += globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = i;

}

fails but

signed int foo = globalReferenceTimeSeries[timeSeriesLength-1];

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  sharedReferenceTimeSeries[i] = i;

}

works (assuming timeSeriesLength>0).

EDIT: Missed your last post. - glad you found the problem already. Unfortunately I don’t know of a tool that would catch such a problem at compile time. Sounds like it would be a fun project to adapt sparse to this.

Ocelot should have been able to catch this at runtime, though.

tera · November 13, 2010, 4:15am

If [font=“Courier New”]foo[/font] is used later, then probably all but the last read of [font=“Courier New”]globalReferenceTimeSeries[i][/font] are optimized away. This would point to a (usually less common) problem at the beginning of the array.

You can verify that theory by checking that

signed int foo = 0;

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  foo += globalReferenceTimeSeries[i];

  sharedReferenceTimeSeries[i] = i;

}

fails but

signed int foo = globalReferenceTimeSeries[timeSeriesLength-1];

for(unsigned int i = 0; i < timeSeriesLength; i++) {

  sharedReferenceTimeSeries[i] = i;

}

works (assuming timeSeriesLength>0).

EDIT: Missed your last post. - glad you found the problem already. Unfortunately I don’t know of a tool that would catch such a problem at compile time. Sounds like it would be a fun project to adapt sparse to this.

Ocelot should have been able to catch this at runtime, though.

Topic		Replies	Views
Unspecified launch failure strange error, please help CUDA Programming and Performance	13	16279	December 31, 2007
unspecified launch failure kernel fails sometimes not everytime CUDA Programming and Performance	6	6069	February 2, 2010
unspecified launch failure kernel fails if a loop is too long CUDA Programming and Performance	8	42842	April 25, 2007
"unspecified launch failure" but "No CUDA-MEMCHECK" CUDA Programming and Performance	7	7519	January 8, 2016
Unspecified launch failure CUDA Programming and Performance	2	5705	May 24, 2009
using cudaMalloc and cudaFree within a loop unspecified launch failure! CUDA Programming and Performance	21	37700	April 23, 2009
Unspecified launch failure on a nested loop (C - multiple pattern matching algorithm) A double loop CUDA Programming and Performance	5	896	November 22, 2010
Unspecifiec launch failure on CUDA_SAFE_CALL(cudaThreadSynchronize()) CUDA Programming and Performance	5	2118	January 27, 2011
A naive question about dereferenced pointers unspecified launch failure CUDA Programming and Performance	10	2927	September 24, 2008
Unspecified launch failure error CUDA Programming and Performance	10	18423	January 6, 2018

YA unspecified launch failure

Related topics