Databases
StackOverflow2013
Postsdbo
POLow performance CUDA code on GT540M
 

Note that there are some explanatory texts on larger screens.

plurals
  1. POLow performance CUDA code on GT540M
    primarykey
    data
    text
    <p>Executing the following code sample takes ~750 ms on a GeForce GT540M whereas the same code executes in ~250 ms on a GT330M.</p> <p>Copying the dev_a and dev_b to the CUDA device memory takes ~350 ms on the GT540M and ~250. The execution of "addCuda" and the copying back to the host takes another ~400 ms on GT540M and ~0 ms for the GT330M.</p> <p>This is not what I expected, so I checked the devices' properties and discovered that the GT540M device surpasses or equals GT330M in every way except in the number of multiprocessors - GT540M has 2 and GT330M has 6. Can this really be true? And if so, can it really have such a great impact on the execution time?</p> <pre><code>#include "cuda_runtime.h" #include "device_launch_parameters.h" #include &lt;stdio.h&gt; #include &lt;stdlib.h&gt; #include &lt;time.h&gt; #include &lt;math.h&gt; #define T 512 #define N 60000*T __global__ void addCuda(double *a, double *b, double *c) { int tid = threadIdx.x + blockIdx.x * blockDim.x; if(tid &lt; N) { c[tid] = sqrt(fabs(a[tid] * b[tid] / 12.34567)) * cos(a[tid]); } } int main() { double *dev_a, *dev_b, *dev_c; double* a = (double*)malloc(N*sizeof(double)); double* b = (double*)malloc(N*sizeof(double)); double* c = (double*)malloc(N*sizeof(double)); printf("Filling arrays (CPU)...\n\n"); int i; for(i = 0; i &lt; N; i++) { a[i] = (double)-i; b[i] = (double)i; } int timer = clock(); cudaMalloc((void**) &amp;dev_a, N*sizeof(double)); cudaMalloc((void**) &amp;dev_b, N*sizeof(double)); cudaMalloc((void**) &amp;dev_c, N*sizeof(double)); cudaMemcpy(dev_a, a, N*sizeof(double), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N*sizeof(double), cudaMemcpyHostToDevice); printf("Memcpy time: %d\n", clock() - timer); addCuda&lt;&lt;&lt;(N+T-1)/T,T&gt;&gt;&gt;(dev_a, dev_b, dev_c); cudaMemcpy(c, dev_c, N*sizeof(double), cudaMemcpyDeviceToHost); printf("Time elapsed: %d\n", clock() - timer); cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); free(a); free(b); free(c); return 0; } </code></pre> <p>The device properties for the devices:</p> <p><strong>GT540M:</strong></p> <pre><code>Major revision number: 2 Minor revision number: 1 Name: GeForce GT 540M Total global memory: 1073741824 Total shared memory per block: 49152 Total registers per block: 32768 Warp size: 32 Maximum memory pitch: 2147483647 Maximum threads per block: 1024 Maximum dimension 0 of block: 1024 Maximum dimension 1 of block: 1024 Maximum dimension 2 of block: 64 Maximum dimension 0 of grid: 65535 Maximum dimension 1 of grid: 65535 Maximum dimension 2 of grid: 65535 Clock rate: 1344000 Total constant memory: 65536 Texture alignment: 512 Concurrent copy and execution: Yes Number of multiprocessors: 2 Kernel execution timeout: Yes </code></pre> <p><strong>GT330M</strong></p> <pre><code>Major revision number: 1 Minor revision number: 2 Name: GeForce GT 330M Total global memory: 268435456 Total shared memory per block: 16384 Total registers per block: 16384 Warp size: 32 Maximum memory pitch: 2147483647 Maximum threads per block: 512 Maximum dimension 0 of block: 512 Maximum dimension 1 of block: 512 Maximum dimension 2 of block: 64 Maximum dimension 0 of grid: 65535 Maximum dimension 1 of grid: 65535 Maximum dimension 2 of grid: 1 Clock rate: 1100000 Total constant memory: 65536 Texture alignment: 256 Concurrent copy and execution: Yes Number of multiprocessors: 6 Kernel execution timeout: Yes </code></pre>
    singulars
    1. PO
      singulars
      1. PTAnswer
    1. This table or related slice is empty.
    1. PTQuestion
    1. This table or related slice is empty.
    1. USThorkil Holm-Jacobsen
    plurals
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PO
      singulars
      1. PTAnswer
    2. PO
      singulars
      1. PTAnswer
    1. VO
      singulars
      1. POLow performance CUDA code on GT540M
      1. This table or related slice is empty.
      1. VTUpMod
    2. VO
      singulars
      1. POLow performance CUDA code on GT540M
      1. This table or related slice is empty.
      1. VTUpMod
    3. VO
      singulars
      1. POLow performance CUDA code on GT540M
      1. This table or related slice is empty.
      1. VTUpMod
    1. COThe first point to make is that the GT330M doesn't support double precision, so you are comparing a single precision result on one device with a double precision one on the other. There is an 8x performance performance difference between the two on current hardware. Can you also add what operating system and CUDA version each is running on to your question, these are mobile parts so they clearly are not running in the same host machine.
      singulars
      1. POLow performance CUDA code on GT540M
      1. UStalonmies
    2. COIn general I think you need to benchmark much larger jobs on CUDA to outweigh the setup / tear-down time. It's surprising the older card can do that more efficiently, though, but is this really a factor in your real application?
      singulars
      1. POLow performance CUDA code on GT540M
      1. USRup
    3. COMachine 1: GT540M (Computing capability 2.1) CUDA version 4.1 Intel Core i5-2410M Windows 7 64bit. Machine 2: GT330M (Computing capability 1.2) CUDA version 4.1 Intel Core i5-520M Windows 7 64bit on bootcamp for Mac. Computing a similar add function on the CPU takes ~2500 ms (on both machines). I tried to replace all doubles with floats to see if it changed anything, but it did not. Does the GPU still use double precision even though it is not required by the application?
      singulars
      1. POLow performance CUDA code on GT540M
      1. USThorkil Holm-Jacobsen
 

Querying!

 
Guidance
A row detail

Detail views are divided into sections. All the information in the data section comes from columns in the selected row. The other sections display data from other, related rows.

Related data can be related in a to-one or a to-many fashion. Captions of data related in a to-many fashion link to a list view showing a filtered view of the table.

Try moving around until you find a non-empty to-many entry and click on the label to get to one. You can move back to the root by clicking on the database name in the header.

SQuiL has stopped working due to an internal error.

If you are curious you may find further information in the browser console, which is accessible through the devtools (F12).

Reload