Opis testu

Celem testu było sprawdzenie jak na czas obliczeń wpływa różny czas życia wątków. W tym celu przyjęto, że wątki które pracują wykonują dokładnie takie same operacje, natomiast czas ich pracy ustalany był na różne sposoby. Sprawdzonych zostało 6 wersji kernela

• Kernel 1:

  __global__ void threadTest_v1(int MAX_THREADS_PER_BLOCK)
  {
      int index = threadIdx.x;
      int long a = 1;
  
  	while(index--){
  		a += a;	
  	}
  }
  

• Kernel 2:

  __global__ void threadTest_v2(int MAX_THREADS_PER_BLOCK)
  {
      int index = threadIdx.x;
  	int iterNumber;
  	int long a = 1;
  
  	if(index % 2 == 0)
  		iterNumber = 1;
  	else
  		iterNumber = MAX_THREADS_PER_BLOCK;
  
  	while(iterNumber--){
  		a += a;
  	}
  }
  

• Kernel 3:

  __global__ void threadTest_v3(int MAX_THREADS_PER_BLOCK)
  {
  	int index = threadIdx.x;
  	int iterNumber;
  	int long a = 1;
  
  	if(index < MAX_THREADS_PER_BLOCK/2)
  		iterNumber = 1;
  	else
  		iterNumber = MAX_THREADS_PER_BLOCK;
  
  	while(iterNumber--){
  		a += a;
  	}
  }
  

• Kernel 4:

  __global__ void threadTest_v4(int MAX_THREADS_PER_BLOCK)
  {
      int index = threadIdx.x;
  	int iterNumber;
  	int long a = 1;
  
  	if(index == 0)
  		iterNumber = 0;
  	else if(index == MAX_THREADS_PER_BLOCK-1)
  		iterNumber = MAX_THREADS_PER_BLOCK;
  	else
  		iterNumber = MAX_THREADS_PER_BLOCK/2;
  
  	while(iterNumber--){
  		a += a;
  	}
  }
  

• Kernel 5:

  __global__ void threadTest_v5(int MAX_THREADS_PER_BLOCK)
  {
  	int iterNumber = MAX_THREADS_PER_BLOCK;
  	int long a = 1;
  
  	while(iterNumber--){
  		a += a;
  	}
  }
  

• Kernel 6:

  __global__ void threadTest_v6(int MAX_THREADS_PER_BLOCK)
  {
  	int index = threadIdx.x;
  	int iterNumber;
  	int long a = 1;
  
  	if(index % 32 == 0)
  		iterNumber = MAX_THREADS_PER_BLOCK;
  	else
  		iterNumber = 1;
  
  	while(iterNumber--){
  		a += a;
  	}
  }
  

Projekt

Download VS2010 project

Wyniki

GeForce 9800 GT

   --- General Information for device 0 ---
Name:  GeForce 9800 GT
Compute capability:  1.1
Clock rate:  1500000
Device copy overlap:  Enabled
Kernel execution timeout :  Enabled
   --- Memory Information for device 0 ---
Total global mem:  536543232
Total constant Mem:  65536
Max mem pitch:  2147483647
Texture Alignment:  256
   --- MP Information for device 0 ---
Multiprocessor count:  14
Shared mem per mp:  16384
Registers per mp:  8192
Threads in warp:  32
Max threads per block:  512
Max thread dimensions:  (512, 512, 64)
Max grid dimensions:  (65535, 65535, 1)

call:
dim3    grids(32768);
dim3    threads(prop.maxThreadsPerBlock);
threadTest_v1<<<grids,threads>>>( prop.maxThreadsPerBlock );

results:
Time for kernel 1:   286.78012 ms
Time for kernel 2:   438.17694 ms
Time for kernel 3:   254.73827 ms
Time for kernel 4:   319.93195 ms
Time for kernel 5:   428.37915 ms
Time for kernel 6:   437.39749 ms

GeForce GTS 450

   --- General Information for device 0 ---
Name:  GeForce GTS 450
Compute capability:  2.1
Clock rate:  1566000
Device copy overlap:  Enabled
Kernel execution timeout :  Enabled
   --- Memory Information for device 0 ---
Total global mem:  536543232
Total constant Mem:  65536
Max mem pitch:  2147483647
Texture Alignment:  512
   --- MP Information for device 0 ---
Multiprocessor count:  4
Shared mem per mp:  49152
Registers per mp:  32768
Threads in warp:  32
Max threads per block:  1024
Max thread dimensions:  (1024, 1024, 64)
Max grid dimensions:  (65535, 65535, 65535)

call:
dim3    grids(32768);
dim3    threads(prop.maxThreadsPerBlock);
threadTest_v1<<<grids,threads>>>( prop.maxThreadsPerBlock );

results:
Time for kernel 1:  1449.91431 ms
Time for kernel 2:  2689.55786 ms
Time for kernel 3:  1184.79163 ms
Time for kernel 4:  1492.06177 ms
Time for kernel 5:  2305.65527 ms
Time for kernel 6:  2732.04346 ms

Wnioski

Porówanie wyników z kernela 2 i 3 oraz 5 i 6 wskazuje na spadek wydajności w sytuacji gdy w jednym bloku występują wątki o różnym czasie życia. W szczególności kernele 2 i 3 wykonują taką samą ilość obliczeń z tym, że kernel 2 z przeplotem co 2 (wątki parzyste wykonują jedną iterację, nieparzyste maksymalną) a kernel 3 grupuje wątki wykonujące po jednej iteracji poczynając od wątku 0 w bloku do MAX_THREADS_PER_BLOCK/2 - 1 oraz maksymalną dla pozostałych.

W przypadku kernela 5 wszystkie wątki wykonują maksymalną ilość iteracji a czas wykonania i tak jest mniejszy od przypadku gdy tylko jeden wątek w warpie wykonuje maksymalną ilość iteracji a pozostałe 31 wykonuje po jednej.

Z niewiadomych przyczyn na GTS 450 kernel 2 powodował niekontrolowane zakończenie programu.