简介

#include <stdio.h>
#include <iostream>
using namespace std;

__global__ void hello_world() {
    printf("Hello world\n");
}

int main(int argc, char** argv) {
    hello_world<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

kernel_name<<<Dg, Db>>>(args);

cudaError_t CUDARTAPI cudaDeviceSynchronize(void);

dim3 gridDim(2,2);
dim3 blockDim(1,2,3);
sumVector<<<gridDim, blockDim>>>(A, B, C);

#include <strings.h>
#include <iostream>
#include <random>
using namespace std;

__global__ void calc(float const* a, float const* b, float* c) {
    unsigned i = threadIdx.x;
    c[i]       = a[i] + b[i];
}

void initData(float* vec, size_t size) {
    static std::random_device             rd;
    static std::uniform_real_distribution<float> dist(0, 1000.f);

    for(int i = 0; i < size; ++i) vec[i] = dist(rd);
}

#define CHECK_CUDA(err)                  \
    {                                        \
        if((err) != cudaSuccess) {           \
            cout << cudaGetErrorString(err); \
            exit(1);                         \
        }                                    \
    }

int main(int argc, char** argv) {
    unsigned size   = 10;
    unsigned nBytes = size * sizeof(float);
    float*   h_A    = new float[size];
    float*   h_B    = new float[size];
    float*   h_C    = new float[size];

    initData(h_A, size);
    initData(h_B, size);
    bzero(h_C, nBytes);
    // 设置 device
    float *d_A, *d_B, *d_C;
    CHECK_CUDA(cudaMalloc((float**)&d_A, nBytes));
    CHECK_CUDA(cudaMalloc((float**)&d_B, nBytes));
    CHECK_CUDA(cudaMalloc((float**)&d_C, nBytes));

    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyKind::cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyKind::cudaMemcpyHostToDevice);

    calc<<<1, 10>>>(d_A, d_B, d_C);

    cudaMemcpy(h_C, d_C, nBytes, cudaMemcpyKind::cudaMemcpyDeviceToHost);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    for(int i = 0; i < size; ++i) {
        std::cout << h_C[i] << ' ';
    }

    cudaDeviceReset();
    return 0;
}

cudaError cudaGetDeviceProperties(cudaDeviceProp *prop, int device)

int main(int argc, char** argv) {
    int  deviceCount;
    auto err = cudaGetDeviceCount(&deviceCount);
    HANDLE_ERROR(err);
    if(!deviceCount) {
        cout << "没有可用于 cuda 的设备" << endl;
        return -1;
    }

    cudaDeviceProp props;
    for(int i = 0; i < deviceCount; ++i) {
        cudaGetDeviceProperties(&props, i);
        cout << "设备名：" << props.name << endl;
        cout << "设备内存存：" << props.totalGlobalMem / 1024 / 1024 << "MB" << endl;
        cout << "设备可用共享内存：" << props.sharedMemPerBlock / 1024 / 1024 << "MB" << endl;
        if(props.integrated) cout << "设备是集成显卡";
    }

    return 0;
}

int getBestDevice() {
    int  deviceCount;
    auto err = cudaGetDeviceCount(&deviceCount);
    HANDLE_ERROR(err);

    int            bestDevice         = -1;
    int            maxMultiProcessors = 0;
    cudaDeviceProp prop;
    for(int i = 0; i < deviceCount; ++i) {
        err = cudaGetDeviceProperties(&prop, i);
        HANDLE_ERROR(err);

        if(prop.multiProcessorCount > maxMultiProcessors) {
            maxMultiProcessors = prop.multiProcessorCount;
            bestDevice         = i;
        }
    }
    return bestDevice;
}

int getGPU() {
    int            device;
    cudaDeviceProp props;

    memset(&props, 0, sizeof(props));
    props.major = 1;
    props.minor = 3;

    cudaError_t err = cudaChooseDevice(&device, &props);
    HANDLE_ERROR(err);
    cout << "设备 " << device << " 的版本号最接近 1.3" << endl;

    return device;
}

float A[M][N];
for (int i = 0; i < M; ++i) {
   for (int j = 0; j < N; ++j) {
        A[i][j] = ...;
   }
}

__global__ void __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) kernel(){}

void __syncthreads();

cudaError_t cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset, cudaMemcpyKind kind);

cudaError_t cudaDeviceSynchronize(void);
__device__ void _syncthreads(void); // 等待线程块中的其它线程达到此点

struct innerStruct {
   float x;
   float yl
}
innerStruct AoS[N];

struct innerArray {
   float x[N];
   float y[N];
}

int tid = blockIdx.x * blockDim.x + threadIdx.x;
int partion = (tid / wrapSize) % 2;
if (partion) {
    //...
}
else {
    //...
}

__global__ void reduceBeughbored(int* g_idata, int* g_odata, unsigned int n) {
   unsigned tid = threadIdx.x;
   int* idata = g_idata + blockIdx.x * blockDim.x;
   int idx = threadIdx.x + blockDim.x * threadIdx.y; // ? 这一行代码对吗？
   if (idx > n) return;

   for (int stride = 1; stride < blockDim.x; stride *= 2) {
      if ((tid % (2 * stride)) == 0) idata[tid] += idata[tid + stride];
      __syncthreads();
   }
   if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

for (int stride = 1; stride < blockDim.x; stride *= 2) {
   int index = 2 * stride * tid;
   if (index < blockDim.x)
      idata[index] += idata[tid + stride];
   __syncthreads();
}

if (tid < 32) {
   volatile int* vmem = idata;
   vmem[tid] += vmem[tid + 32];
   vmem[tid] += vmem[tid + 16];
   vmem[tid] += vmem[tid + 8];
   vmem[tid] += vmem[tid + 4];
   vmem[tid] += vmem[tid + 2];
   vmem[tid] += vmem[tid + 1];
}

// 线程的数量
const size_t THREADS = 5;

/**
 * @param a 输入数据集合
 * @param b 输出数据集合
 * @param size a 的长度，要求为 THREADS 的整数倍
 * b 的长度为 THREADS
 */
__global__ void reduced(int *a, size_t size) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    if(idx > size) return;

    for(int i = THREADS; i < size; i += THREADS) a[idx] = a[idx] + a[idx + i];
}

int computeByCPU(int *arr, size_t size) {
    int res = 0;
    for(int i = 0; i < size; ++i) { res += arr[i]; }
    return res;
}

int main() {
    size_t LENGTH       = THREADS * 200;
    size_t LENGTH_BYTES = LENGTH * sizeof(int);

    int a[LENGTH];
    initArray(a, LENGTH);

    int cpuRes = computeByCPU(a, LENGTH);

    int *d_a;
    auto err = cudaMalloc(&d_a, LENGTH_BYTES);

    err = cudaMemcpy(d_a, a, LENGTH_BYTES, cudaMemcpyHostToDevice);

    reduced<<<1, THREADS>>>(d_a, LENGTH);

    err = cudaDeviceSynchronize();
    err = cudaMemcpy(a, d_a, LENGTH_BYTES, cudaMemcpyDeviceToHost);

    int gpuRes = computeByCPU(a, THREADS);
    cout << "CPU result: " << cpuRes << "\nGPU result: " << gpuRes << endl;

    err = cudaDeviceReset();
    return 0;
}

#include "comm.h"
#include <memory>
#include <omp.h>

const Dim3 gridDim { 12 };
const Dim3 blockDim { 6000000 };

int twoPass(const float* input, size_t size)
{
    float* part_sum = new float[gridDim.x];
    memset(part_sum, 0, sizeof(float) * gridDim.x);
    // 第一步
    // 将 input 中的数据归约到 part_sum
    // [0: size] --> [0: gridDim.x]
    omp_set_num_threads(gridDim.x);
#pragma omp parallel
    {
        int idx = omp_get_thread_num();
        int startIdx = idx * gridDim.x;
        int endIdx = (idx + 1) * gridDim.x;
        endIdx = endIdx > size ? size : endIdx; // 元素不足
        endIdx = idx == (gridDim.x - 1) ? size : endIdx; // 元素多余

        for (; startIdx < endIdx; ++startIdx) {
            part_sum[idx] += input[startIdx];
        }
    }
    // 第二步
    // 将 [0: gridDim.x] --> sum
    for (int i = 1; i < gridDim.x; ++i) {
        part_sum[0] += part_sum[i];
    }
    int res = part_sum[0];
    delete[] part_sum;
    return res;
}

int main(int argc, char** argv)
{
    const size_t DATA_LEN = gridDim.x * blockDim.x + 20;
    float* data = new float[DATA_LEN];
    initArrar(data, DATA_LEN);

    auto start = currentTime();
    int cpuRes = computeSingle(data, DATA_LEN);
    auto point1 = currentTime();
    cout << "cpu 计算消耗的时间为: " << elapsedTime(start, point1) << endl;
    int paralRes = twoPass(data, DATA_LEN);
    delete[] data;
    cout << "parallel 计算消耗的时间为: " << elapsedTime(point1) << endl;
    cout << "cpuRes: " << cpuRes << endl
         << "paralRes: " << paralRes << endl;
    return 0;
}

output[i] = input[0] + input[1] + input[2] + \cdots + input[i]

#pragma omp 指令 子句 子句...

#pragma omp parallel for
for (int i = 0; i < INT_MAX; ++i) {
    result += i;
}

#pragma omp parallel for
for (int i = 0; i < 20; ++i) {
    cout << i << " ";
}

int ReduceSum(int end) {
    int res = 0;
    for (int i = 0; i < end; ++i) {
        res += i;
    }
    return res;
}

#pragma omp parallel for
for (int i = 0; i < 20000; ++i) {
    ReduceSum(i);
}

__global__ void nestedHelloWorld(int iSize, int iDepth) {
   int tid = threadIdx.x;
   printf("Recursion=%d: Hello World from thread %d block %d\n", iDepth, tid, blockIdx.x);
   if (iSize == 1) return;
   int nthreads = iSize >> 1;
   if (tid == 0 && nthreads > 0) {
      nestedHelloWorld <<< 1, nthreads >>> (nthreads, ++iDepth);
      printf("-------> netsted execution depth: %d\n", iDepth);
   }
}