20016. Labeling String (CUDA)

I'm a slow walker, but I never walk backwards.

Problem

現在要求標記一個只包函小寫英文字母和空白的字串 $S$,字串 $S$ 的長度為 $N$ 且每一段連續的英文字母長度 $K$ 至多 500。現在請你用一個 32-bit 整數標記每一個位置上的英文字母前有幾個連續的英文字母 (含當前位置)。例如下述:

S =   a  vy  lu  rah   yhy    fibcl
P = 00100120012001230001230000123450

簡單的循序算法時間複雜度為 $O(N)$。在平行環境下,很容易想到 $O(NK / p + K)$ 的算法,其中 $p$ 為平行處理器的個數。在最簡單的平行算法 $O(NK/p + K)$ 中,每一個執行緒負責一個位置 $P\lbrack i]$ 的結果,往前找到第一個非英文字母的位置 (最多運行 $K$ 次),便可得到所有答案。然而,這樣的方法將會讀取非常多次的 global memory 導致效能變慢。

更進一步地,讓單一執行緒負責一個區塊,讓這個區塊的 $P\lbrack l, r]$ 在線性時間內完成。由於單字長度最大為 $K$,那麼區塊長度必須大於 $K$,便能保證單一執行緒最多運行 $2K $ 次,且讀取 global memory 次數少非常多。然而,這樣的作法在 $K \approx N$ 時退化,又或者 $N/K$ 的個數不足以充分使用整個 GPU 效能。因此,我們可以藉由倍增算法的概念,得到 $O(N \log K / p + \log K)$ 的算法,你可以嘗試使用一些高級數據結構,如稀疏表 (Sparse table)、線段樹 (Segment tree), 二元索引樹 (Binary indexed tree) ... 等。

labeling.h

1
2
3
4
5
6
#ifndef __LABELING_H
#define __LABELING_H
 
void labeling(const char *cuStr, int *cuPos, int strLen);
 
#endif

main.cu

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
#include <bits/stdc++.h>
#include <stdint.h>
#include "labeling.h"
 
#define DEBUG
 
namespace {
    uint32_t seed = 0;
    void p_srand(uint32_t x) { seed = x;}
    uint32_t p_rand() {return seed = (seed*9301 + 49297);}
}
 
int main() {
    cudaSetDevice(1);
    static const int MAXN = 16777216<<1;
    static const int MAXTOKEN = 500;
    static char s[MAXN];
    static int ret[MAXN];
    static char *cuStr;
    static int32_t *cuPos;
    int cases = 0;
    int n, m1, m2, seed;
    while (scanf("%d %d %d %d", &n, &m1, &m2, &seed) == 4) {
        assert(n <= MAXN);
        assert(m1 <= MAXTOKEN);
        p_srand(seed);
        cudaMalloc(&cuStr, sizeof(char)*n);
        cudaMalloc(&cuPos, sizeof(int32_t)*n);
        // random string
        {
            int pos = 0;
            for (; pos < n;) {
                int sp = pos == 0 ? p_rand()%m2 : p_rand()%m2+1;
                int cp = p_rand()%m1+1;
                for (; sp && pos < n; sp--, pos++)
                    s[pos] = ' ';
                for (; cp && pos < n; cp--, pos++)
                    s[pos] = 'a' + p_rand()%26;
            }
#ifdef DEBUG
            for (int i = 0; i < n; i++)
                putchar(s[i]);
            puts("");
#endif
            cudaMemcpy(cuStr, s, sizeof(char)*n, cudaMemcpyHostToDevice);
        }
        // test performance
        {
            clock_t st, ed;
            st = clock();
            const int ROUND = 512;
            for (int i = 0; i < ROUND; i++)    {
                labeling(cuStr, cuPos, n);
            }
            cudaMemcpy(ret, cuPos, sizeof(int32_t)*n, cudaMemcpyDeviceToHost);
            ed = clock() - st;
            fprintf(stderr, "It took average %lf seconds.\n", ((float) ed)/CLOCKS_PER_SEC/ROUND);
        }
        // check
        {
            clock_t st, ed;
            st = clock();
            uint32_t HEX = 0;
            for (int i = 0, sum = 0; i < n; i++) {
#ifdef DEBUG
                // printf("%d%c", ret[i], " \n"[i==n-1]);
                printf("%d", ret[i]);
                if (i == n-1)
                    puts("");
#endif
                if (s[i] > ' ')
                    sum++;
                else
                    sum = 0;
                assert(sum == ret[i]);
                HEX ^= i*ret[i];
            }
            ed = clock() - st;
            fprintf(stderr, "Check task took %lf seconds.\n", ((float) ed)/CLOCKS_PER_SEC);
            printf("Case #%d: PASS %X\n", ++cases, HEX);
        }
        cudaFree(cuPos);
        cudaFree(cuStr);
    }
 
    return 0;
}

Input Format

輸入有多組測資,每組測資一行包含四個整數 $N, \; M1, \; M2, \text{seed}$,自動產生字串長度 $N$,連續字母長度最大 $M1$、連續空白長度最大 $M2$ 以及亂數種子 $\text{seed}$。

  • $1 \le N \le 10^7$
  • $1 \le M1, \; M2 \le 500 $

Output Format

根據上述程序會驗證並輸出雜湊值,若答案輸出錯誤則會發生 $\texttt{Wrong Answer}$ 的訊息。

Sample Input

32 5 4 1

Sample Output (stdout)

Case #1: PASS B4

Sample Output (stderr)

  a  vy  lu  rah   yhy    fibcl
It took average 0.000011 seconds.
00100120012001230001230000123450
Check task took 0.000006 seconds.

Compile

all: main
 
main: main.cu labeling.cu
        nvcc -arch sm_30 -Xcompiler "-O2 -fopenmp" -c labeling.cu -o labeling.o
        nvcc -arch sm_30 -Xcompiler "-O2 -fopenmp" main.cu labeling.o -o main

Source

GPU Programming 2017 Spring, Author: Yu Sheng Lin, Instructor: Wei Chao Chen

Discussion