diff options
| -rw-r--r-- | expt_0516.cu | 72 |
1 files changed, 28 insertions, 44 deletions
diff --git a/expt_0516.cu b/expt_0516.cu index 351a8c4..1954380 100644 --- a/expt_0516.cu +++ b/expt_0516.cu @@ -8,6 +8,9 @@ #include <cstdio> #include <vector> +constexpr auto NUM_THREADS = 512; +constexpr auto FACTOR = 1; + std::vector<unsigned long long> population_vector, sample_vector, le_sample_vector; @@ -54,70 +57,57 @@ inline void cudaCheckError(const char *file, const int line) { } */ -__global__ void kernel(unsigned long step, const double slice_size) { +__global__ void kernel(const size_t test_length) { auto custom_sort = CustomSort(cuda_sample_length, sizeof(key_type) * 8); // printf("kernel1 step: %ld\n", step); - for (int i = 0; i < step; i++) { - auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; + for (int i = 0; i < FACTOR; i++) { + auto tid = + i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= test_length) + return; // printf("%d\n", tid); - auto _index = (int)(custom_sort.sample_cdf_custom_version( - cudaSampleItem, cudaPopulationItem[tid]) / - slice_size) - - 1; + custom_sort.sample_cdf_custom_version(cudaSampleItem, + cudaPopulationItem[tid]); // cdf_result[index] = true; } } -__global__ void kernel2(unsigned long step, const double slice_size) { +__global__ void kernel2(const size_t test_length) { - for (int i = 0; i < step; i++) { - auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i; - auto _index = - (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length, - cudaPopulationItem[tid]) / - slice_size) - - 1; - // cdf_result[index] = true; - } -} - -__global__ void print2() { - /*for (int i = 0; i< SAMPLE_LENGTH; i++) { - printf("%llu ", cudaSampleItem[i]); - } - printf("\n");*/ - printf("%lu\n", CustomSort::fast_log(cuda_sample_length)); -} + for (int i = 0; i < FACTOR; i++) { + auto tid = + i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x; -__global__ void printFunction() {} + if (tid >= test_length) + return; -__global__ void checkItems() { - for (int i = 0; i < cuda_sample_length; i++) { - printf("%lld ", cudaSampleItem[i]); + FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length, + cudaPopulationItem[tid]); + // cdf_result[index] = true; } - printf("\n"); } -constexpr dim3 grid_dim = 32, block_dim = 32; +inline size_t calcBlocks(size_t input) { + return (input % (NUM_THREADS * FACTOR) == 0) + ? input / (NUM_THREADS * FACTOR) + : (input / (NUM_THREADS * FACTOR)) + 1; +} void run_kernel(size_t test_size, bool custom = true) { // need description - unsigned long step = test_size / (grid_dim.x * block_dim.x); // Only scale mode will use - auto scale = 2; - const unsigned long split_size = test_size * scale; - const auto slice_size = 1.0 / (double)split_size; + auto blocks = calcBlocks(test_size); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, nullptr); if (custom) { - kernel<<<grid_dim, block_dim>>>(step, slice_size); + kernel<<<blocks, NUM_THREADS>>>(test_size); } else { - kernel2<<<grid_dim, block_dim>>>(step, slice_size); + kernel2<<<blocks, NUM_THREADS>>>(test_size); } cudaDeviceSynchronize(); @@ -132,9 +122,6 @@ void run_kernel(size_t test_size, bool custom = true) { } printf("%stime: %lf ", custom ? "custom " : "", time); cudaDeviceSynchronize(); - - printFunction<<<1, 1>>>(); - cudaDeviceSynchronize(); } __global__ void freeStorageStage1() { cudaFree(cudaSampleItem); } @@ -229,9 +216,6 @@ int main(int argc, char const *argv[]) { "kernel\n"); run_kernel(test_size); - // freeStorageStage1<<<1, 1>>>(); - // cudaDeviceSynchronize(); - cudaFree(cudaSample); key_type *cudaNormalSample = nullptr; |