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" notifications@github.com Gesendet: ‎04.‎02.‎2014 19:07 An: "cbuchner1/CudaMiner" CudaMiner@noreply.github.com 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 notifications@github.com:

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 .