Skip to content

Commit

Permalink
cuda: Support CUDA toolkit versions older than 10.
Browse files Browse the repository at this point in the history
This commit adds a conditional macro to the Blake3 CUDA kernel so that
on older CUDA Toolkit versions a shift-based rotation is used instead of
the intrinsic.

According to the CUDA Toolkit documentation[1], the __funnelshift_rc
integer intrinsic was introduced on CUDA version 10. Prior to this
commit, the CUDA kernel of gominer could not be built when using
Toolkit versions older than 10.

Note that although this makes it possible to build the kenel on older
Toolkit versions, performance is significantly degraded when compared to
the modern, intrinsic-based versions.

[1]: https://docs.nvidia.com/cuda/archive/10.0/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html
  • Loading branch information
matheusd authored and davecgh committed Sep 18, 2023
1 parent 9f1c3f0 commit 1c92e04
Showing 1 changed file with 5 additions and 0 deletions.
5 changes: 5 additions & 0 deletions blake3.cu
Expand Up @@ -15,7 +15,12 @@
#define MAX_OUTPUT_RESULTS 32

// Written and optimized by Dave Collins Sep 2023.
#if __CUDACC_VER_MAJOR__ >= 10
#define ROTR(v, n) __funnelshift_rc((v), (v), n)
#else
#define ROTR(v, n) ((v) >> n) | ((v) << (32 - n))
#endif


__global__
void search(
Expand Down

0 comments on commit 1c92e04

Please sign in to comment.