thrust

#include <stdio.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <iostream>
using namespace std;

__host__ __device__
int sort_func(int a, int b){
    return a > b;
}

int main(){

    int data[] = {5, 3, 1, 5, 2, 0};
    int ndata  = sizeof(data) / sizeof(data[0]);
    thrust::host_vector<int> array1(data, data + ndata);
    thrust::sort(array1.begin(), array1.end(), sort_func);

    thrust::device_vector<int> array2 = thrust::host_vector<int>(data, data + ndata);
    thrust::sort(array2.begin(), array2.end(), []__device__(int a, int b){return a < b;});

    printf("array1------------------------\n");
    for(int i = 0; i < array1.size(); ++i)
        cout << array1[i] << endl;

    printf("array2------------------------\n");
    for(int i = 0; i < array2.size(); ++i)
        cout << array2[i] << endl;
    return 0;
}

cuda错误机制

# 知识点
1. 若cuda核函数出错，由于他是异步的，立即执行cudaPeekAtLastError只会拿到对输入参数校验是否正确的状态，而不会拿到核函数是否执行正确的状态
2. 因此需要等待核函数执行完毕后，才真的知道当前核函数是否出错，一般通过设备同步或者流同步进行等待
3. 错误分为可恢复和不可恢复两种：
    - 可恢复：
        - 参数配置错误等，例如block越界（一般最大值是1024），shared memory大小超出范围（一般是48KB）
        - 通过cudaGetLastError可以获取错误代码，同时把当前状态恢复为success
        - 该错误在调用核函数后可以立即通过cudaGetLastError/cudaPeekAtLastError拿到
        - 该错误在下一个函数调用的时候会覆盖
    - 不可恢复：
        - 核函数执行错误，例如访问越界等等异常
        - 该错误则会传递到之后的所有cuda操作上
        - 错误状态通常需要等到核函数执行完毕才能够拿到，也就是有可能在后续的任何流程中突然异常（因为是异步的）

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

__global__ void func(float* ptr){

    int pos = blockIdx.x * blockDim.x + threadIdx.x;
    if(pos == 999){
        ptr[999] = 5;
    }
}

int main(){

    float* ptr = nullptr;

    // 因为核函数是异步的，因此不会立即检查到他是否存在异常
    func<<<100, 10>>>(ptr);
    //func<<<100, 1050>>>(ptr);
    auto code1 = cudaPeekAtLastError();
    cout << cudaGetErrorString(code1) << endl;

    // 对当前设备的核函数进行同步，等待执行完毕，可以发现过程是否存在异常
    auto code2 = cudaDeviceSynchronize();
    cout << cudaGetErrorString(code2) << endl;

    // 异常会一直存在，以至于后续的函数都会失败
    float* new_ptr = nullptr;
    auto code3 = cudaMalloc(&new_ptr, 100);
    cout << cudaGetErrorString(code3) << endl;
    return 0;
}