CUDA addition vs shift instruction performance
===

I am trying to understand instruction throughput for a large CUDA kernel that I am working on. I wrote two small programs to compare the throughput of addition and shift instructions. According to CUDA C Programming Guide, the throughput for shift instruction is half of add instruction. However, when I measure the time of following two programs on Tesla M2070, the time is exactly same. Can someone please explain why this is the case?

Addition Program:

    #include <limits.h>
    #include <stdio.h>
    #include <fstream>
    #include <iostream>
    #include <cstdlib>
    #include <stdint.h>
    
    using namespace std;
    
    __global__ void testAdd(int numIterations, uint1* result){
      int total = 1;
      for(int i=0; i< numIterations;i ++){
        total = total+i;
      }
      result[0] = make_uint1(total);
    }
    
    int main(){
      uint1* result;
      cudaMalloc((void**)(&(result)), sizeof(uint1));
      float totalElapsedTime = 0;
      int i;
    
      for(i = 0; i < 10; i++){
        cudaEvent_t start, stop;
        cudaEventCreate(&start);
        cudaEventCreate(&stop);
        cudaEventRecord(start, 0);
        testAdd<<<1,1>>>(100000, result);
        cudaError_t e50 = cudaGetLastError();
        if(e50 == cudaSuccess){
          cudaEventRecord(stop, 0);
          cudaEventSynchronize(stop);
          float elapsedTime;
          cudaEventElapsedTime(&elapsedTime, start, stop);
          totalElapsedTime += elapsedTime;;
          //cout << "Elapsed Time:" << elapsedTime << endl;                                                                                                                              
    
        }else{
          cout << "Error launching kernel: " << e50 << endl;
        }
        cudaEventDestroy(start);
        cudaEventDestroy(stop);
      }
      cout << "Elapsed Time: " << totalElapsedTime/i << endl;
      cudaFree(result);
    }

Shift Program:

    #include <limits.h>
    #include <stdio.h>
    #include <fstream>
    #include <iostream>
    #include <cstdlib>
    #include <stdint.h>
    
    using namespace std;
    
    __global__ void testShift(int numIterations, uint1* result){
      int total = 1;
      for(int i=0; i< numIterations;i ++){
        total = total<<i;
      }
      result[0] = make_uint1(total);
    }
    
    int main(){
      uint1* result;
      cudaMalloc((void**)(&(result)), sizeof(uint1));
      float totalElapsedTime = 0;
      int i;
    
      for(i = 0; i < 10; i++){
        cudaEvent_t start, stop;
        cudaEventCreate(&start);
        cudaEventCreate(&stop);
        cudaEventRecord(start, 0);
        testShift<<<1,1>>>(100000, result);
        cudaError_t e50 = cudaGetLastError();
        if(e50 == cudaSuccess){
          cudaEventRecord(stop, 0);
          cudaEventSynchronize(stop);
          float elapsedTime;
          cudaEventElapsedTime(&elapsedTime, start, stop);
          totalElapsedTime += elapsedTime;;
          //cout << "Elapsed Time:" << elapsedTime << endl;                                                                                                                              
    
        }else{
          cout << "Error launching kernel: " << e50 << endl;
        }
        cudaEventDestroy(start);
        cudaEventDestroy(stop);
      }
      cout << "Elapsed Time: " << totalElapsedTime/i << endl;
      cudaFree(result);
    }