title: cudaStreamWaitEvent does not seem to wait
body: <p>I am attempting to write a small demo program that has two cuda streams progressing and, governed by events, waiting for each other. So far this program looks like this:</p>
  
  <pre><code>// event.cu
  #include &lt;iostream&gt;
  #include &lt;cstdio&gt;
  #include &lt;cuda_runtime.h&gt;
  #include &lt;cuda.h&gt;
  
  using namespace std;
  
  __global__ void k_A1() { printf("\tHi! I am Kernel A1.\n"); }
  __global__ void k_B1() { printf("\tHi! I am Kernel B1.\n"); }
  __global__ void k_A2() { printf("\tHi! I am Kernel A2.\n"); }
  __global__ void k_B2() { printf("\tHi! I am Kernel B2.\n"); }
  
  int main()
  {
    cudaStream_t streamA, streamB;
    cudaEvent_t halfA, halfB;
    cudaStreamCreate(&amp;streamA);
    cudaStreamCreate(&amp;streamB);
    cudaEventCreate(&amp;halfA);
    cudaEventCreate(&amp;halfB);
  
    cout &lt;&lt; "Here is the plan:" &lt;&lt; endl &lt;&lt;
      "Stream A: A1, launch 'HalfA', wait for 'HalfB', A2." &lt;&lt; endl &lt;&lt;
      "Stream B: Wait for 'HalfA', B1, launch 'HalfB', B2." &lt;&lt; endl &lt;&lt;
      "I would expect: A1,B1, (A2 and B2 running concurrently)." &lt;&lt; endl;
  
    k_A1&lt;&lt;&lt;1,1,0,streamA&gt;&gt;&gt;(); // A1!
    cudaEventRecord(halfA,streamA); // StreamA triggers halfA!
    cudaStreamWaitEvent(streamA,halfB,0); // StreamA waits for halfB.
    k_A2&lt;&lt;&lt;1,1,0,streamA&gt;&gt;&gt;(); // A2!
  
    cudaStreamWaitEvent(streamB,halfA,0); // StreamB waits, for halfA.
    k_B1&lt;&lt;&lt;1,1,0,streamB&gt;&gt;&gt;(); // B1!
    cudaEventRecord(halfB,streamB); // StreamB triggers halfB!
    k_B2&lt;&lt;&lt;1,1,0,streamB&gt;&gt;&gt;(); // B2!
  
    cudaEventDestroy(halfB);
    cudaEventDestroy(halfA);
    cudaStreamDestroy(streamB);
    cudaStreamDestroy(streamA);
  
    cout &lt;&lt; "All has been started. Synchronize!" &lt;&lt; endl;
    cudaDeviceSynchronize();
    return 0;
  }
  </code></pre>
  
  <p>My grasp of CUDA streams is the following: A stream is a kind of list to which I can add tasks. These tasks are tackled in series. So in my program I can rest assured that streamA would in order</p>
  
  <ol>
  <li>Call kernel k_A1</li>
  <li>Trigger halfA</li>
  <li>Wait for someone to trigger halfB</li>
  <li>Call kernel k_A2</li>
  </ol>
  
  <p>and streamB would</p>
  
  <ol>
  <li>Wait for someone to trigger halfA</li>
  <li>Call kernel k_B1</li>
  <li>Trigger halfB</li>
  <li>Call kernel k_B2</li>
  </ol>
  
  <p>Normally both streams might run asynchronous to each other. However, I would like to block streamB until A1 is done and then block streamA until B1 is done.</p>
  
  <p>This appears not to be as simple. On my Ubuntu with Tesla M2090 (CC 2.0) the output of</p>
  
  <pre><code>nvcc -arch=sm_20 event.cu &amp;&amp; ./a.out
  </code></pre>
  
  <p>is</p>
  
  <pre><code>Here is the plan:
  Stream A: A1, launch 'HalfA', wait for 'HalfB', A2.
  Stream B: Wait for 'HalfA', B1, launch 'HalfB', B2.
  I would expect: A1,B1, (A2 and B2 running concurrently).
  All has been started. Synchronize!
          Hi! I am Kernel A1.
          Hi! I am Kernel A2.
          Hi! I am Kernel B1.
          Hi! I am Kernel B2.
  </code></pre>
  
  <p>And I really would have expected B1 to be completed before the cudaEventRecord(halfB,streamB). Nevertheless stream A obviously does not wait for the completion of B1 and so not for the recording of halfB.</p>
  
  <p>What's more: If I altogether delete the cudaEventRecord commands I would expect the program to lock down on the cudaStreamWait commands. But it does not and produces the same output. What am I overlooking here?</p>