2016-06-28

批改娘 10099. Dot Product (CUDA)

1. 題目描述
1. 1.1. main.c
2. 1.2. utils.h
2. 範例輸入
3. 範例輸出
4. 編譯參數
5. Solution

題目描述

請用 CUDA 改寫下段的計算：

main.c

#include <stdio.h>
#include <assert.h>
#include <omp.h>
#include <inttypes.h>
#include <stdint.h>
#include "utils.h"
 
#define MAXGPU 8
#define MAXCODESZ 32767
#define MAXN 16777216
uint32_t A[MAXN], B[MAXN], C[MAXN];
int main(int argc, char *argv[]) {
    omp_set_num_threads(4);
    int N;
    uint32_t key1, key2;
    while (scanf("%d %" PRIu32 " %" PRIu32, &N, &key1, &key2) == 3) {
        int chunk = N / 4;
        for (int i = 0; i < N; i++) {
            A[i] = encrypt(i, key1);
            B[i] = encrypt(i, key2);
        }
 
        for (int i = 0; i < N; i++)
            C[i] = A[i] * B[i];
 
        uint32_t sum = 0;
        for (int i = 0; i < N; i++)
            sum += C[i];
        printf("%" PRIu32 "\n", sum);
    }
    return 0;
}

utils.h

#ifndef _UTILS_H
#define _UTILS_H
#include <stdint.h>
static inline uint32_t rotate_left(uint32_t x, uint32_t n) {
    return  (x << n) | (x >> (32-n));
}
static inline uint32_t encrypt(uint32_t m, uint32_t key) {
    return (rotate_left(m, key&31) + key)^key;
}
#endif

範例輸入

1 2	16777216 1 2 16777216 3 5

範例輸出

1 2	2885681152 2147483648

編譯參數

1 2	$ nvcc -Xcompiler "-O2 -fopenmp" main.cu -o main $ ./main

Solution

這裡同我們在 OpenCL 的實作技巧，將生成測資和計算都丟在 GPU 上完成，但是 CUDA 只能在 Nvidia 顯卡上運作，而且根據版本的不同，每一種顯卡的計算能力也不同，可以參考 wiki，最低版本為 1.0，也就在編譯參數中加入 nvcc -arch=compute_10，如果可以到 2.0，下達 nvcc -arch=compute_20，以此類推。編譯器預設計算能力為 1.0，因此如果要在 kernel function 裡面印出訊息 (意即 printf())，至少提供 2.0 以上的編譯參數。

CUDA 程式撰寫就不用像 OpenCL 從找尋 Platform 到抓到 Device，之後再藉由 Device IDs 建立 Context，再從 Context 建立 Program，CUDA 提供 特殊語法，而不像 OpenCL 採用 特殊函數 包裝，這導致編程複雜度差異極大，但是從彈性來看 OpenCL 可以調控的項目較多且動態，但 CUDA 由於是自家產品，效能會稍微比同版本的 OpenCL 來得快，一部分也是因為編譯器不同導致的緣故。

#include <stdio.h>
#include <stdint.h>
#include <cuda.h>
#include <omp.h>
__device__ uint32_t rotate_left(uint32_t x, uint32_t n) {
    return  (x << n) | (x >> (32-n));
}
__device__ uint32_t encrypt(uint32_t m, uint32_t key) {
    return (rotate_left(m, key&31) + key)^key;
}
__host__ uint32_t h_rotate_left(uint32_t x, uint32_t n) {
    return  (x << n) | (x >> (32-n));
}
__host__ uint32_t h_encrypt(uint32_t m, uint32_t key) {
    return (h_rotate_left(m, key&31) + key)^key;
}
#define MAXN 16777216
#define GPULOCAL 128
#define BLOCKSZ (1024)
__global__ void vecdot(uint32_t keyA, uint32_t keyB, uint32_t C[], int N) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int l = x * BLOCKSZ;
    int r = l + BLOCKSZ;
    uint32_t sum = 0;
    if (r > N)	r = N;
    for (int i = l; i < r; i++)
        sum += encrypt(i, keyA) * encrypt(i, keyB);
    C[x] = sum;
}
uint32_t hostC[MAXN / GPULOCAL];
#define CheckErr(status) { gpuAssert((status), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, int abort=true) {
   if (code != cudaSuccess) {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}
int main() {
    uint32_t N, keyA, keyB;
    uint32_t *cuArrC;
    cudaMalloc((void **)&cuArrC, MAXN/GPULOCAL*sizeof(uint32_t));
    while (scanf("%u %u %u", &N, &keyA, &keyB) == 3) {
        int M = (N + BLOCKSZ-1) / BLOCKSZ;
        int LOCAL = GPULOCAL;
        M = (M + LOCAL) / LOCAL * LOCAL;
        dim3 cuBlock(LOCAL);
        dim3 cuGrid(M/LOCAL);
        vecdot<<<cuGrid, cuBlock>>>(keyA, keyB, cuArrC, N);
        CheckErr(cudaGetLastError());
        cudaMemcpy(hostC, cuArrC, M*sizeof(uint32_t), cudaMemcpyDeviceToHost);
        uint32_t sum = 0;
#ifdef _OPENMP
        omp_set_num_threads(4);
#endif
        #pragma omp parallel for reduction(+: sum)
        for (int i = 0; i < M; i++)
            sum += hostC[i];
        printf("%u\n", sum);
    }
    cudaFree(cuArrC);
    return 0;
}

Morris' Blog