Function Calls are not handled correctly in the emulator or parser

[GoogleCodeExporter]

What steps will reproduce the problem?
1. Compile any CUDA device function with the directive __noinline__

What is the expected output? What do you see instead?

Before we were under the impression that the compiler never generated any
function calls.  It turns out that it does.  So we should provide support
for this in the parser and the emulator.

We need to add unit tests for the lexer/parser that contain device function
calls. We also need to add functionality to the code generator for the
emulator to add code of all referenced function calls to a kernel binary. 
Finally, we need to add support for recursive function calls since they are
likely to come out in future releases...

Original issue reported on code.google.com by gregory....@gatech.edu on 7 Jul 2009 at 2:08

[GoogleCodeExporter]

This is an example program that causes nvcc to generate function calls:

#define SIZE 100

__noinline__ __device__ float value( ) {
  return 0.0f;
}

__global__ void kernel(float* x) {
  const unsigned p = blockIdx.x*blockDim.x + threadIdx.x;
  x[p] = value();
}

int main(int argc, char** argv) {
  float* x;
  cudaMalloc((void**)&x, SIZE*sizeof(float));

  dim3 Dg, Db;
  Db.x = SIZE;

  kernel<<<Dg,Db>>>(x);

  return 0;
}

The resulting source and PTX files are attached.  These should be added as unit 
tests.

Original comment by gregory....@gatech.edu on 13 Feb 2010 at 8:10

Attachments:

• deadlock.ptx
• deadlock.cu

[GoogleCodeExporter]

This is implemented in the PTX 2.1 branch

Original comment by gregory....@gatech.edu on 18 Oct 2010 at 6:47

• Changed state: Fixed