петли разворачивая с динамическим параллелизмом снизить производительность времени

Question

У меня есть простая программа для вычисления квадратного корня, размотка цикл был сделана как

петли разворачивая

#include <stdio.h> 
#include <cuda.h> 
__global__ void square(float *a, int N,int idx); 


// Kernel that executes on the CUDA device 
__global__ void first(float *arr, int N) 
{ 
    int idx = 2*(blockIdx.x * blockDim.x + threadIdx.x); 
    int n=N; 
    //printf("%d\n",n); 
    for(int q=0;q<2;q++) 
    { 
    if(N<2000) 
    { 
    arr[idx+q] = arr[idx+q] * arr[idx+q]; 
    } 
    } 

} 



// main routine that executes on the host 
int main(void) 
{ 
    clock_t start = clock(),diff; 
    float *a_h, *a_d; // Pointer to host & device arrays 
    const int N = 1000; // Number of elements in arrays 
    size_t size = N * sizeof(float); 
    a_h = (float *)malloc(size);  // Allocate array on host 
    cudaMalloc((void **) &a_d, size); // Allocate array on device 
    // Initialize host array and copy it to CUDA device 
    for (int i=0; i<N; i++) a_h[i] = (float)i; 
    cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); 
    // Do calculation on device: 
    int block_size = 4; 
    //int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); 
    first <<< 4, 128 >>> (a_d, N); 
    //cudaThreadSynchronize(); 
    // Retrieve result from device and store it in host array 
    cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); 
    // Print results 
    for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]); 
    // Cleanup 
    free(a_h); cudaFree(a_d); 
    diff = clock() - start; 
int msec = diff * 1000/CLOCKS_PER_SEC; 

printf("Time taken %d seconds %d milliseconds\n", msec/1000, msec%1000); 

} 

затем понимая, что вычисление контура может быть сведены к минимуму динамический параллелизм.

разворачивая с динамическим параллелизмом был реализован как

разворачивая с динамическим параллелизмом

#include <stdio.h> 
#include <cuda.h> 
__global__ void square(float *a, int N,int idx); 


// Kernel that executes on the CUDA device 
__global__ void first(float *arr, int N) 
{ 
    int idx = 2*(blockIdx.x * blockDim.x + threadIdx.x); 
    int n=N; 
    square <<< 1,2 >>> (arr, n,idx); 


} 

__global__ void square(float *a, int N,int idx) 
{ 
    int tdx = blockIdx.x * blockDim.x + threadIdx.x; 
    printf("%d\n",N); 
    if(N<2000) 
    { 
    a[tdx+idx] = a[tdx+idx] * a[tdx+idx]; 
    } 
} 

// main routine that executes on the host 
int main(void) 
{ 
    clock_t start = clock(),diff; 
    float *a_h, *a_d; // Pointer to host & device arrays 
    const int N = 1000; // Number of elements in arrays 
    size_t size = N * sizeof(float); 
    a_h = (float *)malloc(size);  // Allocate array on host 
    cudaMalloc((void **) &a_d, size); // Allocate array on device 
    // Initialize host array and copy it to CUDA device 
    for (int i=0; i<N; i++) a_h[i] = (float)i; 
    cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); 
    // Do calculation on device: 
    int block_size = 4; 
    //int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); 
    first <<< 4, 128 >>> (a_d, N); 
    //cudaThreadSynchronize(); 
    // Retrieve result from device and store it in host array 
    cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); 
    // Print results 
    for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]); 
    // Cleanup 
    free(a_h); cudaFree(a_d); 
    diff = clock() - start; 
int msec = diff * 1000/CLOCKS_PER_SEC; 

printf("Time taken %d seconds %d milliseconds\n", msec/1000, msec%1000); 

} 

реализация динамического параллелизма разворачивания занимает больше времени, чем executio только разворачивая. Арен, мы предполагаем улучшить время выполнения с динамическим параллелизмом в таком случае?

Answer 1

Динамический параллелизм в основном полезен в случаях, когда у вас есть параллелизм, который является динамическим. То есть: случаи, когда вы не знаете, сколько параллелизма вам понадобится, пока вы не выполните некоторые вычисления. Вместо того, чтобы передавать данные обратно на хост, который затем мгновенно подается в параметризацию другого запуска, вы запускаете изнутри ядра. В этом шаблоне, когда memcpys между запущенными ядрами запускается, вы увидите ускорение.

В приведенном выше примере это не тот случай. Вы могли бы запустить только два потока от хоста. Нет ничего динамического, поскольку нет параллелизма, доступного там, о котором вы не знали во время первого запуска ядра.

Кроме того, требования к производительности для ядер, запущенных с использованием динамического параллелизма, аналогичны требованиям производительности, запущенным с хоста. Вы должны запустить разумный объем работы, или латентность запуска будет доминировать над временем вычислений.