Lecture 10: Parallel Merge Assignment

Scaffold Head

#include <iostream> template <class value_t, class index_t> static __device__ __host__ __forceinline__ index_t get_rank_inclusive(value_t val, value_t * data, index_t L) { if (L == 0) return 0; index_t pos = 0, stride = L; for (; stride > 0; stride >>= 1) { index_t newPos = pos + stride < L ? pos + stride : L; if (data[newPos - 1] <= val) pos = newPos; } return pos; } template <class value_t, class index_t> static __device__ __host__ __forceinline__ index_t get_rank_exclusive(value_t val, value_t * data, index_t L) { if (L == 0) return 0; index_t pos = 0, stride = L; for (; stride > 0; stride >>= 1) { index_t newPos = pos + stride < L ? pos + stride : L; if (data[newPos - 1] < val) pos = newPos; } return pos; }

Scaffold Foot

#include <vector> int main () { // from the lecture std::vector<int> output(16, 0); std::vector<int> input = {5, 9, 10, 10, 17, 19, 40, 50, 3, 6, 10, 10, 21, 25, 45, 55}; std::vector<int> sb_len_left = {0, 4, 2, 2}; std::vector<int> sb_len_right = {1, 3, 0, 4}; // initialize memory int * Input, * Output, * Sb_len_left, * Sb_len_right; cudaMalloc(&Input, sizeof(int)*16); cudaMalloc(&Output, sizeof(int)*16); cudaMalloc(&Sb_len_left, sizeof(int)*4); cudaMalloc(&Sb_len_right, sizeof(int)*4); // copy memory auto H2D = cudaMemcpyHostToDevice; cudaMemset(Output, 0, sizeof(int)*16); cudaMemcpy(Input, input.data(), sizeof(int)*16, H2D); cudaMemcpy(Sb_len_left, sb_len_left.data(), sizeof(int)*4, H2D); cudaMemcpy(Sb_len_right, sb_len_right.data(), sizeof(int)*4, H2D); for (int sb = 0; sb < 4; sb++) { // incrementally call kernel to visualize the algorithm pairwise_merge<<<1, 4>>>(Input, Output, Sb_len_left, Sb_len_right, sb+1, 8); cudaMemcpy(output.data(), Output, sizeof(int)*16, cudaMemcpyDeviceToHost); // print output to command line after phase sb std::cout << "processed subblock number " << sb << std::endl; for (const auto& value : output) std::cout << value << " "; std::cout << std::endl; } cudaFree(Input); cudaFree(Output); cudaFree(Sb_len_left); cudaFree(Sb_len_right); }

Start time:

Do 12 Apr 2018 16:11:00

End time:

Fr 01 Mär 2019 12:00:00

General test timeout:

10.0 seconds

Tests

processed subblock number 0
3 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 
processed subblock number 1
3 5 6 9 10 10 10 10 0 0 0 0 0 0 0 0 
processed subblock number 2
3 5 6 9 10 10 10 10 17 19 0 0 0 0 0 0 
processed subblock number 3
3 5 6 9 10 10 10 10 17 19 21 25 40 45 50 55