Error checks on grid launches?

[jtynefield]

Hi Christian,

As I play around with the parameter space, some configurations will fail to launch, but the failures are not caught and are then expressed as crazy performance and failed work. Any reason not to add error checking after all the grid launch directives?

Even checking it only once would catch most of these.

....
        switch(opt_algo) {
          case ALGO_SCRYPT:      fermi_scrypt_core_kernelA_LG<ALGO_SCRYPT><<< grid, threads, shared, stream >>>(d_idata, N, LOOKUP_GAP); break;
          case ALGO_SCRYPT_JANE: fermi_scrypt_core_kernelA_LG<ALGO_SCRYPT_JANE><<< grid, threads, shared, stream >>>(d_idata, N, LOOKUP_GAP); break;
        }

        static int once = 1;
        if ( once ) {
           once = 0;
           cudaError_t error = cudaGetLastError();
           if ( error != cudaSuccess ) {
              fprintf( stderr, "Error: %s: %s\n", "grid launch", cudaGetErrorString(error) );
              exit( -1 );
           }
        }

[cbuchner1]

Yeah I am pretty aware of these shortcomings in my error checking

----- Ursprüngliche Nachricht ----- Von: "jtynefield" [email protected] Gesendet: ‎04.‎02.‎2014 19:07 An: "cbuchner1/CudaMiner" [email protected] Betreff: [CudaMiner] Error checks on grid launches? (#88)

Hi Christian, As I play around with the parameter space, some configurations will fail to launch, but the failures are not caught and are then expressed as crazy performance and failed work. Any reason not to add error checking after all the grid launch directives? Even checking it only once would catch most of these. .... switch(opt_algo) { case ALGO_SCRYPT: fermi_scrypt_core_kernelA_LG<<< grid, threads, shared, stream >>>(d_idata, N, LOOKUP_GAP); break; case ALGO_SCRYPT_JANE: fermi_scrypt_core_kernelA_LG<<< grid, threads, shared, stream >>>(d_idata, N, LOOKUP_GAP); break; } static int once = 1; if ( once ) { once = 0; cudaError_t error = cudaGetLastError(); if ( error != cudaSuccess ) { fprintf( stderr, "Error: %s: %s\n", "grid launch", cudaGetErrorString(error) ); exit( -1 ); } } — Reply to this email directly or view it on GitHub.

[jtynefield]

If I coded them up, would you accept the diffs?

[cbuchner1]

if you are familiar with git, you can create a branch and submit a pull request with your changes. I could try to integrate it then.

2014-02-04 jtynefield [email protected]:

If I coded them up, would you accept the diffs?

Reply to this email directly or view it on GitHubhttps://github.com/cbuchner1/CudaMiner/issues/88#issuecomment-34108120 .