2016-10-03 2 views
1

다른 커널 구현을 비교하기 위해 clock()을 사용하고 싶습니다. 간단한 SAXPY 예제에서 구현하려고 시도했지만 제로 클록주기가 발생하지만 이는 거의 없습니다.CUDA clock()은 제로 클럭 사이클을 유도합니다.

이미 clock()을 구현하는 방법에 대한 몇 가지 예제를 발견했습니다. herehere. 하지만 어떻게 든 내 코드로의 전송이 작동하지 않습니다.

/* SAXPY code example from https://devblogs.nvidia.com/parallelforall/easy-introduction-cuda-c-and-c/ */ 

#include <stdio.h> 

// The declaration specifier __global__ defines a kernel. This code 
// will be copied to the device and will be executed there in parallel 
__global__ 
void saxpy(int n, float a, float *x, float *y, int *kernel_clock) 
{ 
    // The indexing of the single threads is done with the following 
    // code line 
    int i = blockIdx.x*blockDim.x + threadIdx.x; 

    clock_t start = clock(); 

    // Each thread is executing just one position of the arrays 
    if (i < n) y[i] = a*x[i] + y[i]; 

    clock_t stop = clock(); 

    kernel_clock[i] = (int) (stop-start); 
} 

int main(void) 
{ 
    // Clock cycles of threads 
    int *kernel_clock; 
    int *d_kernel_clock; 
    // Creating a huge number 
    int N = 1<<20; 
    float *x, *y, *d_x, *d_y; 
    // Allocate an array on the *host* of the size of N 
    x = (float*)malloc(N*sizeof(float)); 
    y = (float*)malloc(N*sizeof(float)); 
    kernel_clock = (int*)malloc(N*sizeof(int)); 

    // Allocate an array on the *device* of the size of N 
    cudaMalloc(&d_x, N*sizeof(float)); 
    cudaMalloc(&d_y, N*sizeof(float)); 
    cudaMalloc(&d_kernel_clock, N*sizeof(int)); 

    // Filling the array of the host 
    for (int i = 0; i < N; i++) { 
    x[i] = 1.0f; 
    y[i] = 2.0f; 
    } 

    // Copy the host array to the device array 
    cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_kernel_clock, kernel_clock, N*sizeof(int), cudaMemcpyHostToDevice); 

    // Perform SAXPY on 1M elements. The triple chevrons dedicates how 
    // the threads are grouped on the device 
    saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y, d_kernel_clock); 
    cudaDeviceSynchronize(); 

    // Copy the result from the device to the host 
    cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost); 
    cudaMemcpy(kernel_clock, d_kernel_clock, N*sizeof(int), cudaMemcpyDeviceToHost); 

    // Calculate average clock time 
    float average_clock = 0; 
    for (int i = 0; i < N; i++) { 
     average_clock += (float) (kernel_clock[i]); 
    } 
    average_clock /= N; 

    // Display the time to the screen 
    printf ("Kernel clock cycles: %.4f\n", average_clock); 

    // Free the memory on the host and device 
    free(x); 
    free(y); 
    free(kernel_clock); 
    cudaFree(d_x); 
    cudaFree(d_y); 
    cudaFree(d_kernel_clock); 
} 

이 코드 예제는 리드 (Lead) : 내가 잘못하고있는 무슨 확실하지 않다

Kernel clock cycles: 0.0000 

여기

내가 사용하고있는 코드입니다. 그래서 제 질문은 : 실제로 합리적인 결과를 얻으려면 어떻게해야합니까? 당신이 당신의 질문에 링크 된 답변 중 하나에서 인용

+0

오류 검사가 표시되지 않습니다. 'cuda-memcheck'로 코드를 실행하면 어떻게됩니까? –

+0

'cuda-memcheck'는 0 에러를냅니다. '======== 에러 요약 : 0 errors' – stebran

답변

1

또한 컴파일러 및 어셈블러가 검사 할 수 있도록 명령 재정렬을 수행 할 수 있음을 알고 있어야합니다

그 시계 전화 SASS 출력 (cuobjdump를 사용하여 확인)에서 서로 옆에 놓는 것을 피하십시오.

나는 이것이 귀하의 문제의 근원이라고 생각합니다. 나는 CUDA 8 릴리스 툴킷을 사용하여 커널을 컴파일 한 후 cuobjdump와 결과 기계 코드를 분해하면, 나는 다음과 같은 얻을 :

code for sm_52 
      Function : _Z5saxpyifPfS_Pi 
    .headerflags @"EF_CUDA_SM52 EF_CUDA_PTX_SM(EF_CUDA_SM52)" 
                          /* 0x001c4400fe0007f6 */ 
    /*0008*/     MOV R1, c[0x0][0x20];          /* 0x4c98078000870001 */ 
    /*0010*/   {   CS2R R7, SR_CLOCKLO;          /* 0x50c8000005070007 */ 
    /*0018*/     S2R R0, SR_CTAID.X;  }        /* 0xf0c8000002570000 */ 
                          /* 0x083fc400e3e007f0 */ 
    /*0028*/   {   CS2R R8, SR_CLOCKLO;          /* 0x50c8000005070008 */ 
    /*0030*/     S2R R2, SR_TID.X;  }         /* 0xf0c8000002170002 */ 
    /*0038*/     XMAD.MRG R3, R0.reuse, c[0x0] [0x8].H1, RZ;     /* 0x4f107f8000270003 */ 
                          /* 0x081fc400fec207f6 */ 
    /*0048*/     XMAD R2, R0.reuse, c[0x0] [0x8], R2;      /* 0x4e00010000270002 */ 
    /*0050*/     XMAD.PSL.CBCC R0, R0.H1, R3.H1, R2;       /* 0x5b30011800370000 */ 
    /*0058*/     ISETP.GE.AND P0, PT, R0.reuse, c[0x0][0x140], PT;   /* 0x4b6d038005070007 */ 
                          /* 0x001fd400fc2007ec */ 
    /*0068*/     SHR R9, R0, 0x1f;           /* 0x3829000001f70009 */ 
    /*0070*/    @!P0 SHF.L.U64 R2, RZ, 0x2, R0;         /* 0x36f800400028ff02 */ 
    /*0078*/    @!P0 SHF.L.U64 R3, R0, 0x2, R9;         /* 0x36f804c000280003 */ 
                          /* 0x001fc040fe4207f6 */ 
    /*0088*/    @!P0 IADD R4.CC, R2.reuse, c[0x0][0x148];      /* 0x4c10800005280204 */ 
    /*0090*/    @!P0 IADD.X R5, R3.reuse, c[0x0][0x14c];       /* 0x4c10080005380305 */ 
    /*0098*/   { @!P0 IADD R2.CC, R2, c[0x0][0x150];        /* 0x4c10800005480202 */ 
    /*00a8*/    @!P0 LDG.E R4, [R4];  }         /* 0x0005c400fe400076 */ 
                          /* 0xeed4200000080404 */ 
    /*00b0*/    @!P0 IADD.X R3, R3, c[0x0][0x154];        /* 0x4c10080005580303 */ 
    /*00b8*/    @!P0 LDG.E R6, [R2];            /* 0xeed4200000080206 */ 
                          /* 0x001fd800fea007e1 */ 
    /*00c8*/     LEA R10.CC, R0, c[0x0][0x158], 0x2;       /* 0x4bd781000567000a */ 
    /*00d0*/     IADD R8, -R7, R8;           /* 0x5c12000000870708 */ 
    /*00d8*/     LEA.HI.X R9, R0, c[0x0][0x15c], R9, 0x2;     /* 0x1a17048005770009 */ 
                          /* 0x001fc008fe4007f1 */ 
    /*00e8*/     MOV R7, R9;             /* 0x5c98078000970007 */ 
    /*00f0*/    @!P0 FFMA R0, R4, c[0x0][0x144], R6;        /* 0x4980030005180400 */ 
    /*00f8*/   {   MOV R6, R10;            /* 0x5c98078000a70006 */ 
    /*0108*/    @!P0 STG.E [R2], R0;  }         /* 0x001ffc005e2001f2 */ 
                          /* 0xeedc200000080200 */ 
    /*0110*/     STG.E [R6], R8;            /* 0xeedc200000070608 */ 
    /*0118*/     EXIT;              /* 0xe30000000007000f */ 
                          /* 0x001f8000fc0007ff */ 
    /*0128*/     BRA 0x120;             /* 0xe2400fffff07000f */ 
    /*0130*/     NOP;              /* 0x50b0000000070f00 */ 
    /*0138*/     NOP;              /* 0x50b0000000070f00 */ 
      ................................. 

당신은 그들이 어떤없이 호출되도록 클럭 지침이 다시 정렬 된 것을 볼 수 있습니다 그 (것)들의 사이에서 암호로하십시오. 이 코드를 실행하는 모든 워프가 아닌 모든 경우에 대해 제로 또는 매우 제로 클럭 측정에 가깝습니다.

+0

고마워요! 이제는 문제를 이해하지만 출력물에서 어떤 선이 표시됩니까? – stebran

관련 문제