asyncAPI

标签：__ const int asyncAPI nbytes data checkCudaErrors

async.cu

#include <stdio.h>

#include <cuda_runtime.h>
#include <cuda_profiler_api.h>

template <typename T>
void check(T result, char const *const func, const char *const file,
            int const line)
{
    if (result) {
        fprintf(stderr, "CUDA error at %s:%d code=%d(%s) \"%s\" \n", file, line,
               static_cast<unsigned int>(result), cudaGetErrorString(result), func);
        exit(EXIT_FAILURE);
    }
}

#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)

__global__ void increment_kernel(int *g_data, int inc_value)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    g_data[idx] = g_data[idx] + inc_value;
}

bool correct_output(int *data, const int n, const int x)
{
    for (int i = 0; i < n; i++) {
        if (data[i] != x) {
            printf("Error! data[%d] = %d, ref = %d\n", i, data[i], x);
            return false;
        }
    }

    return true;
}


int main(int argc, char *argv[])
{
    printf("[%s] - Starting...\n", argv[0]);

    checkCudaErrors(cudaSetDevice(0));

    int n = 16 * 1024 * 1024;
    int nbytes = n * sizeof(int);
    int value = 26;

    int *a = 0;
    checkCudaErrors(cudaMallocHost((void **)&a, nbytes));
    memset(a, 0, nbytes);

    int *d_a = 0;
    checkCudaErrors(cudaMalloc((void **)&d_a, nbytes));
    checkCudaErrors(cudaMemset(d_a, 255, nbytes));

    dim3 threads = dim3(512, 1);
    dim3 blocks = dim3(n / threads.x, 1);

    cudaEvent_t start, stop;
    checkCudaErrors(cudaEventCreate(&start));
    checkCudaErrors(cudaEventCreate(&stop));

    checkCudaErrors(cudaDeviceSynchronize());
    float gpu_time = 0.0f;

    checkCudaErrors(cudaProfilerStart());
    cudaEventRecord(start, 0);
    cudaMemcpyAsync(d_a, a, nbytes, cudaMemcpyHostToDevice, 0);
    increment_kernel<<<blocks, threads, 0, 0>>>(d_a, value);
    cudaMemcpyAsync(a, d_a, nbytes, cudaMemcpyDeviceToHost, 0);
    cudaEventRecord(stop, 0);
    checkCudaErrors(cudaProfilerStop());

    unsigned long int counter = 0;

    while (cudaEventQuery(stop) == cudaErrorNotReady) {
        counter++;
    }

    checkCudaErrors(cudaEventElapsedTime(&gpu_time, start, stop));

    printf("time spent executting by GPU: %.2f\n", gpu_time);
    printf("CPU executed %lu iterations while waiting for GPU to finish\n", counter);

    bool bFinalResults = correct_output(a, n, value);

    checkCudaErrors(cudaEventDestroy(start));
    checkCudaErrors(cudaEventDestroy(stop));
    checkCudaErrors(cudaFreeHost(a));
    checkCudaErrors(cudaFree(d_a));

    exit(bFinalResults ? EXIT_SUCCESS : EXIT_FAILURE);
}

运行结果：

标签：__,const,int,asyncAPI,nbytes,data,checkCudaErrors
From： https://www.cnblogs.com/ZhangYouJiu/p/18686949