2

我在使用 CUDA 时遇到以下性能问题。当我在 Titan V 和 Titan X 卡上运行一个简单的示例代码时，运行时间符合预期。

Titan X: 0.269299 ms
Titan V: 0.111766 ms

现在，当我在代码中添加另一个内核时，它使用动态并行，但仍然不调用它或根本不使用它，Volta GPU 的性能急剧下降，但在其他卡上性能不受影响。

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
4

0 回答 0