Powyższe odpowiedzi wskazują, jak rozmiar bloku może wpływać na wydajność i sugerują typową heurystykę jej wyboru opartą na maksymalizacji obłożenia. Nie chcąc zapewnić kryterium wyboru rozmiaru bloku, warto byłoby wspomnieć, że CUDA 6.5 (obecnie w wersji Release Candidate) zawiera kilka nowych funkcji, aby pomóc w czasie wykonywania obliczeń obłożenia i konfiguracji startu, patrz
Wskazówka CUDA Pro: API Occupancy upraszcza konfigurację uruchamiania
Jedną z przydatnych funkcji jest cudaOccupancyMaxPotentialBlockSizeheurystyczna kalkulacja rozmiaru bloku, który zapewnia maksymalną liczbę osób. Wartości dostarczane przez tę funkcję można następnie wykorzystać jako punkt wyjścia do ręcznej optymalizacji parametrów uruchamiania. Poniżej znajduje się mały przykład.
#include <stdio.h>
/************************/
/* TEST KERNEL FUNCTION */
/************************/
__global__ void MyKernel(int *a, int *b, int *c, int N)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < N) { c[idx] = a[idx] + b[idx]; }
}
/********/
/* MAIN */
/********/
void main()
{
const int N = 1000000;
int blockSize; // The launch configurator returned block size
int minGridSize; // The minimum grid size needed to achieve the maximum occupancy for a full device launch
int gridSize; // The actual grid size needed, based on input size
int* h_vec1 = (int*) malloc(N*sizeof(int));
int* h_vec2 = (int*) malloc(N*sizeof(int));
int* h_vec3 = (int*) malloc(N*sizeof(int));
int* h_vec4 = (int*) malloc(N*sizeof(int));
int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int));
int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int));
int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int));
for (int i=0; i<N; i++) {
h_vec1[i] = 10;
h_vec2[i] = 20;
h_vec4[i] = h_vec1[i] + h_vec2[i];
}
cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice);
float time;
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N);
// Round up according to array size
gridSize = (N + blockSize - 1) / blockSize;
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("Occupancy calculator elapsed time: %3.3f ms \n", time);
cudaEventRecord(start, 0);
MyKernel<<<gridSize, blockSize>>>(d_vec1, d_vec2, d_vec3, N);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
printf("Kernel elapsed time: %3.3f ms \n", time);
printf("Blocksize %i\n", blockSize);
cudaMemcpy(h_vec3, d_vec3, N*sizeof(int), cudaMemcpyDeviceToHost);
for (int i=0; i<N; i++) {
if (h_vec3[i] != h_vec4[i]) { printf("Error at i = %i! Host = %i; Device = %i\n", i, h_vec4[i], h_vec3[i]); return; };
}
printf("Test passed\n");
}
EDYTOWAĆ
Element cudaOccupancyMaxPotentialBlockSizejest zdefiniowany w cuda_runtime.hpliku i jest zdefiniowany w następujący sposób:
template<class T>
__inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize(
int *minGridSize,
int *blockSize,
T func,
size_t dynamicSMemSize = 0,
int blockSizeLimit = 0)
{
return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit);
}
Znaczenie parametrów jest następujące
minGridSize = Suggested min grid size to achieve a full machine launch.
blockSize = Suggested block size to achieve maximum occupancy.
func = Kernel function.
dynamicSMemSize = Size of dynamically allocated shared memory. Of course, it is known at runtime before any kernel launch. The size of the statically allocated shared memory is not needed as it is inferred by the properties of func.
blockSizeLimit = Maximum size for each block. In the case of 1D kernels, it can coincide with the number of input elements.
Zauważ, że od CUDA 6.5, trzeba obliczyć własne wymiary bloku 2D / 3D z rozmiaru bloku 1D sugerowanego przez API.
Należy również zauważyć, że interfejs API sterownika CUDA zawiera funkcjonalnie równoważne interfejsy API do obliczania zajętości, więc możliwe jest użycie cuOccupancyMaxPotentialBlockSizew kodzie interfejsu API sterownika w taki sam sposób, jak pokazano w przypadku interfejsu API środowiska wykonawczego w powyższym przykładzie.