diff options
| author | KunoiSayami <[email protected]> | 2023-05-18 03:40:26 +0800 |
|---|---|---|
| committer | KunoiSayami <[email protected]> | 2023-05-18 03:40:26 +0800 |
| commit | d38b99aba596b1a4edbc7e296f826d5ad936d357 (patch) | |
| tree | d74a7d00c223f11b25adeed4db05e7f8345839e6 | |
| parent | 1ad1e0ebf58da30ee7f65a3a1f2a8f62900c47d5 (diff) | |
fix(exp): Add missing function in expt_0516_2
Signed-off-by: KunoiSayami <[email protected]>
| -rw-r--r-- | expt_0516.cu | 4 | ||||
| -rw-r--r-- | expt_0516_2.cu | 26 | ||||
| -rw-r--r-- | sortlib.cuh | 8 |
3 files changed, 24 insertions, 14 deletions
diff --git a/expt_0516.cu b/expt_0516.cu index f141027..82755ac 100644 --- a/expt_0516.cu +++ b/expt_0516.cu @@ -260,12 +260,12 @@ int main(int argc, char const *argv[]) { // custom_sort.testCalculation(); P_ERR("\rApply custom length and test calculation"); applyCudaSampleLength<<<1, 1>>>(sample_length); - testCustomCalculation<<<1, 1>>>(sample_length, sizeof(key_type) * 8); + // testCustomCalculation<<<1, 1>>>(sample_length); cudaDeviceSynchronize(); P_ERR("\rApply custom length and test calculation completed"); // Can remove this function if pass - CustomSort::testSelf(sample_length, sizeof(key_type) * 8); + // CustomSort::testSelf(sample_length); initSample<<<1, 1>>>(cudaSample, cudaPopulation); cudaDeviceSynchronize(); diff --git a/expt_0516_2.cu b/expt_0516_2.cu index 0661e47..860e146 100644 --- a/expt_0516_2.cu +++ b/expt_0516_2.cu @@ -11,37 +11,47 @@ class CustomSort { public: explicit CustomSort(size_t length) : LENGTH(length) {} const size_t LENGTH; - static size_t fast_log(size_t a) { - float t = a; - return (((*(int *)&t) >> 23) + 1) & 127; + + __device__ __host__ static size_t fast_log(size_t a) { +#ifdef __CUDA_ARCH__ + return (size_t)log2((double)a); +#else + return (size_t)std::log2(a); +#endif } - size_t calculate_index(size_t rank) const { + __device__ __host__ size_t calculate_index(size_t rank) const { size_t bit_low = (LENGTH + 1) >> fast_log(++rank) >> 1; return (((rank << 1) | 1) * bit_low - LENGTH - 1); } - size_t calculate_rank(size_t index) const { + __device__ __host__ size_t calculate_rank(size_t index) const { index++; size_t low_bit = index & (-index); return ((LENGTH + index) / low_bit) >> 1; } - __attribute__((unused)) void testCalculation() const { + __device__ __host__ void testCalculation() const { for (size_t i = 0; i < LENGTH; i++) { auto left = calculate_index(i); auto right = calculate_rank(left - 1); - if (i + 1 == LENGTH || i + 1 != right) { + /*if (i + 1 == LENGTH || i + 1 != right) { printf("%lu %lu %lu\n", i, left, right); printf("%lu %lu %lu\n", LENGTH, fast_log(i + 1), (LENGTH + 1) >> fast_log(i + 1) >> 1); - } + }*/ assert(i + 1 == right); } } static void testSelf(size_t length) { CustomSort(length).testCalculation(); } }; + +__global__ void testCustomCalculationWithPrint(size_t length) { + CustomSort(length).testCalculation(); + printf("%lu pass\n", length); +} + #endif long pow_for_sample(long n) { diff --git a/sortlib.cuh b/sortlib.cuh index 5d32e4b..5bacc11 100644 --- a/sortlib.cuh +++ b/sortlib.cuh @@ -19,9 +19,9 @@ public: __device__ __host__ static size_t fast_log(size_t a) { #ifdef __CUDA_ARCH__ - return log2((double)a); + return (size_t)log2((double)a); #else - return std::log2(a); + return (size_t)std::log2(a); #endif } @@ -40,11 +40,11 @@ public: for (size_t i = 0; i < LENGTH; i++) { auto left = calculate_index(i); auto right = calculate_rank(left - 1); - if (i + 1 == LENGTH || i + 1 != right) { + /*if (i + 1 == LENGTH || i + 1 != right) { printf("%lu %lu %lu\n", i, left, right); printf("%lu %lu %lu\n", LENGTH, fast_log(i + 1), (LENGTH + 1) >> fast_log(i + 1) >> 1); - } + }*/ assert(i + 1 == right); } } |
