Снижение производительности в Titan V при наличии кода динамического параллелизма » - PullRequest
Новая
       55

Снижение производительности в Titan V при наличии кода динамического параллелизма »

0 голосов
/ 11 октября 2018

У меня возникла следующая проблема с производительностью в CUDA.Когда я запускаю простой пример кода на картах Titan V и Titan X, время выполнения отлично, как и ожидалось. 

Titan X: 0.269299 ms
Titan V: 0.111766 ms

Теперь, когда я добавляю в код другое ядро, которое использует динамический параллелизм, новсе равно не вызывайте его или не используйте его вообще, производительность в графическом процессоре Volta резко падает, но на других картах производительность не влияет. 

Titan X: 0.270602 ms
Titan V: 1.999299 ms

Важно подчеркнуть тот факт, что этот второйЯдро вообще не используется, оно просто находится рядом с остальным кодом, т. е. оно компилируется только с остальным кодом.Можно также прокомментировать рекурсивные вызовы ядра вместе с созданием потока, и увидеть, что время работы Volta снова стало хорошим.Я подозреваю, что наличие динамического параллелизма негативно влияет на код, даже если он вообще не используется во время выполнения.Любые идеи о том, как подойти к этой проблеме? 

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <math.h>
__global__ void MetNormal(int *a_d,const int N );
__global__ void rec(int *a_d ,int Nm,int xi, int yi, int zi, const int BlockSize){

    int x=blockIdx.x*blockDim.x+threadIdx.x+xi;
    int y=blockIdx.y*blockDim.y+threadIdx.y+yi;
    int z=blockIdx.z*blockDim.z+threadIdx.z+zi;

    int Nbloques= (Nm+BlockSize-1)/BlockSize;
    dim3 b(BlockSize,BlockSize,BlockSize);
    dim3 g(Nbloques,Nbloques,Nbloques);

    cudaStream_t s1;//line of code to comment
    cudaStreamCreateWithFlags(&s1,cudaStreamNonBlocking);// line of code to comment
    rec<<<g,b,0,s1>>>(a_d,Nm,xi,yi,zi,BlockSize);//line of code to comment

}
__global__ void MetNormal(int *a_d,const int N){

    int x= blockIdx.x*blockDim.x+threadIdx.x;
    int y= blockIdx.y*blockDim.y+threadIdx.y;
    int z= blockIdx.z*blockDim.z+threadIdx.z;
    int ind=z*N*N+y*N+x;

   a_d[ind]=1;
}
int main(int argc ,char **argv){

    if (argc !=4){
        fprintf(stderr,"Error, run program as ./prog N rep device\n");
        exit(EXIT_FAILURE);
    }

    unsigned long N=atoi(argv[1]);
    unsigned long rep=atoi(argv[2]);
    cudaSetDevice(atoi(argv[3]));
    int *a,*a_d, xi=0, yi=0,zi=0;
    int BSize=8;

    a=(int*)malloc(sizeof(int)*N*N*N);
    cudaMalloc((void ** ) &a_d,N*N*N*sizeof(int));

    dim3 Bloque(BSize,BSize,BSize);
    float NB=(float)N/(float)(2*BSize);
    int B=(int) ceil(NB);

    dim3 GridBruto((N+BSize-1)/BSize,(N+BSize-1)/BSize,(N+BSize-1)/BSize);

    fflush(stdout);
    for(int i=0;i<N;i++){
        for (int j=0;j<N;j++){
            for(int k=0;k<N;k++){
                a[N*N*k+i*N+j]=0;
            }
        }
    }

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaMemcpy(a_d,a,N*N*N*sizeof(int),cudaMemcpyHostToDevice); 

    cudaEventRecord(start);
    for(int i =0;i<rep;i++){
        MetNormal<<<GridBruto,Bloque>>>(a_d,N); 
        cudaDeviceSynchronize();
        }
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);               
    printf("Time %f ms\n", milliseconds/(rep));
    fflush(stdout);

    cudaDeviceSynchronize();
    cudaMemcpy(a,a_d,N*N*N*sizeof(int),cudaMemcpyDeviceToHost);
    return 0;
}

строка компиляции: 

nvcc -O3 -std=c++11 -lm -arch sm_60 -rdc=true -lcudadevrt prog.cu -o prog
...