summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2023-05-18 03:40:26 +0800
committerKunoiSayami <[email protected]>2023-05-18 03:40:26 +0800
commitd38b99aba596b1a4edbc7e296f826d5ad936d357 (patch)
treed74a7d00c223f11b25adeed4db05e7f8345839e6
parent1ad1e0ebf58da30ee7f65a3a1f2a8f62900c47d5 (diff)
fix(exp): Add missing function in expt_0516_2
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--expt_0516.cu4
-rw-r--r--expt_0516_2.cu26
-rw-r--r--sortlib.cuh8
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);
}
}