github-jwinterm/xmr-stak
1
0
Fork 0
mirror of https://github.com/jwinterm/xmr-stak.git synced 2026-10-20 06:08:37 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #399 from psychocrypt/topic-nvidiaErrorWithMessage

Browse Source 
add message to `CUDA_CHECK...` macros
This commit is contained in:
fireice-uk
2017-12-08 11:04:50 +00:00
committed by GitHub
parent 99b681453a e5127b5ffa
commit 3bc72ec9c0
3 changed files with 53 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files
xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
+16 -12
View File
@@ -327,18 +327,22 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx)
for ( int i = 0; i < partcount; i++ )
{
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<ITERATIONS,THREAD_SHIFT,MASK><<<
grid,
block4,
block4.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 )
>>>(
ctx->device_blocks*ctx->device_threads,
ctx->device_bfactor,
i,
ctx->d_long_state,
ctx->d_ctx_a,
ctx->d_ctx_b
));
CUDA_CHECK_MSG_KERNEL(
ctx->device_id,
"\n**suggestion: Try to increase the value of the attribute 'bfactor' or \nreduce 'threads' in the NVIDIA config file.**",
cryptonight_core_gpu_phase2<ITERATIONS,THREAD_SHIFT,MASK><<<
grid,
block4,
block4.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 )
>>>(
ctx->device_blocks*ctx->device_threads,
ctx->device_bfactor,
i,
ctx->d_long_state,
ctx->d_ctx_a,
ctx->d_ctx_b
)
);
if ( partcount > 1 && ctx->device_bsleep > 0) compat_usleep( ctx->device_bsleep );
}
xmrstak/backend/nvidia/nvcc_code/cuda_device.hpp
+28 -9
View File
@@ -9,22 +9,41 @@
/** execute and check a CUDA api command
*
* @param id gpu id (thread id)
* @param msg message string which should be added to the error message
* @param ... CUDA api command
*/
#define CUDA_CHECK(id, ...) { \
cudaError_t error = __VA_ARGS__; \
if(error!=cudaSuccess){ \
std::cerr << "[CUDA] Error gpu " << id << ": <" << __FILE__ << ">:" << __LINE__ << std::endl; \
throw std::runtime_error(std::string("[CUDA] Error: ") + std::string(cudaGetErrorString(error))); \
} \
} \
#define CUDA_CHECK_MSG(id, msg, ...) { \
cudaError_t error = __VA_ARGS__; \
if(error!=cudaSuccess){ \
std::cerr << "[CUDA] Error gpu " << id << ": <" << __FILE__ << ">:" << __LINE__; \
std::cerr << msg << std::endl; \
throw std::runtime_error(std::string("[CUDA] Error: ") + std::string(cudaGetErrorString(error))); \
} \
} \
( (void) 0 )
/** execute and check a CUDA api command
*
* @param id gpu id (thread id)
* @param ... CUDA api command
*/
#define CUDA_CHECK(id, ...) CUDA_CHECK_MSG(id, "", __VA_ARGS__)
/** execute and check a CUDA kernel
*
* @param id gpu id (thread id)
* @param ... CUDA kernel call
*/
#define CUDA_CHECK_KERNEL(id, ...) \
__VA_ARGS__; \
#define CUDA_CHECK_KERNEL(id, ...) \
__VA_ARGS__; \
CUDA_CHECK(id, cudaGetLastError())
/** execute and check a CUDA kernel
*
* @param id gpu id (thread id)
* @param msg message string which should be added to the error message
* @param ... CUDA kernel call
*/
#define CUDA_CHECK_MSG_KERNEL(id, msg, ...) \
__VA_ARGS__; \
CUDA_CHECK_MSG(id, msg, cudaGetLastError())
xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu
+9 -2
View File
@@ -218,7 +218,6 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
}
size_t wsize = ctx->device_blocks * ctx->device_threads;
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, hashMemSize * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key1, 40 * sizeof(uint32_t) * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key2, 40 * sizeof(uint32_t) * wsize));
@@ -228,6 +227,10 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_input, 21 * sizeof (uint32_t ) ));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_result_count, sizeof (uint32_t ) ));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_result_nonce, 10 * sizeof (uint32_t ) ));
CUDA_CHECK_MSG(
ctx->device_id,
"\n**suggestion: Try to reduce the value of the attribute 'threads' in the NVIDIA config file.**",
cudaMalloc(&ctx->d_long_state, hashMemSize * wsize));
return 1;
}
@@ -254,7 +257,11 @@ extern "C" void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce,
CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_nonce, 0xFF, 10 * sizeof (uint32_t ) ));
CUDA_CHECK(ctx->device_id, cudaMemset( ctx->d_result_count, 0, sizeof (uint32_t ) ));
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_final<<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state ));
CUDA_CHECK_MSG_KERNEL(
ctx->device_id,
"\n**suggestion: Try to increase the value of the attribute 'bfactor' in the NVIDIA config file.**",
cryptonight_extra_gpu_final<<<grid, block >>>( wsize, target, ctx->d_result_count, ctx->d_result_nonce, ctx->d_ctx_state )
);
CUDA_CHECK(ctx->device_id, cudaMemcpy( rescount, ctx->d_result_count, sizeof (uint32_t ), cudaMemcpyDeviceToHost ));
CUDA_CHECK(ctx->device_id, cudaMemcpy( resnonce, ctx->d_result_nonce, 10 * sizeof (uint32_t ), cudaMemcpyDeviceToHost ));