0_Simple__cppIntegration

▶ 分离编译【留坑，在 Linux 上用命令行试一下】

▶ 源代码：

// cppIntegration_gold.cpp
#include <vector_types.h>

extern "C" void computeGold(char *reference, char *idata, const unsigned int len);
extern "C" void computeGold2(int2 *reference, int2 *idata, const unsigned int len);

void computeGold(char *reference, char *idata, const unsigned int len)
{
    for (unsigned int i = 0; i < len; ++i)
        reference[i] = idata[i] - 10;
}

void computeGold2(int2 *reference, int2 *idata, const unsigned int len)
{
    for (unsigned int i = 0; i < len; ++i)
    {
        reference[i].x = idata[i].x - idata[i].y;
        reference[i].y = idata[i].y;
    }
}

// cppIntegration.cu
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <helper_cuda.h>
#include <helper_functions.h>

#ifndef MAX
#define MAX(a,b) (a > b ? a : b)
#endif

extern "C" void computeGold(char *reference, char *idata, const unsigned int len);
extern "C" void computeGold2(int2 *reference, int2 *idata, const unsigned int len);

// GPU上的运算
__global__ void kernel(int *g_data)
{
    const unsigned int tid = threadIdx.x;
    int data = g_data[tid];
    // data 每个字节的值减去 10，再拼接到一起
    g_data[tid] = ((((data <<  0) >> 24) - 10) << 24) | ((((data <<  8) >> 24) - 10) << 16) | ((((data << 16) >> 24) - 10) <<  8) | ((((data << 24) >> 24) - 10) <<  0) ;
}

__global__ void kernel2(int2 *g_data)       // 使用 int2* 格式的输入
{
    const unsigned int tid = threadIdx.x;
    int2 data = g_data[tid];    
    g_data[tid].x = data.x - data.y;        // data.x 中每个元素减去 data.y 中对应元素的偏移量
}

// 测试不同的核函数处理的结果。输入两种格式的待处理数据，及其长度
extern "C" bool runTest(char *data, int2 *data_int2, unsigned int len)
{   
    assert((len % 4) == 0);                                                     // 要求数组长度为 4 的倍数
    const unsigned int num_threads = len / 4, mem_size = sizeof(char) * len, mem_size_int2 = sizeof(int2) * len;

    char *d_data;
    cudaMalloc((void **)&d_data, mem_size);
    cudaMemcpy(d_data, data, mem_size, cudaMemcpyHostToDevice);
    int2 *d_data_int2;
    cudaMalloc((void **)&d_data_int2, mem_size_int2);
    cudaMemcpy(d_data_int2, data_int2, mem_size_int2, cudaMemcpyHostToDevice);
    
    kernel << < dim3(1, 1, 1), dim3(num_threads, 1, 1) >> > ((int *)d_data);
    kernel2 << < dim3(1, 1, 1), dim3(len, 1, 1) >> > (d_data_int2);

    getLastCudaError("Kernel execution failed");                                // 检查和函数运行是否有错误，有错则输出这话

    char *reference = (char *)malloc(mem_size);                                 // 使用 CPU 计算
    computeGold(reference, data, len);
    printf("ref char*:%s
", reference);
    int2 *reference2 = (int2 *)malloc(mem_size_int2);
    computeGold2(reference2, data_int2, len);
    printf("ref int2 :");
    for (int i = 0; i < len;i++)
        printf("%c", reference2[i].x);
    printf("
");

    cudaMemcpy(data, d_data, mem_size, cudaMemcpyDeviceToHost);    
    cudaMemcpy(data_int2, d_data_int2, mem_size_int2, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    printf("gpu char*:%s
", (char *)data);
    printf("gpu int2 :");
    for (int i = 0; i < len; i++)        
        printf("%c", data_int2[i].x);
    printf("
");
    
    cudaFree(d_data);
    cudaFree(d_data_int2);
    free(reference);
    free(reference2);
    return 0;
}

// main.cpp
#include <iostream>
#include <cstdlib>
#include <cuda_runtime.h>
#include <vector_types.h>
#include <helper_cuda.h>

extern "C" bool runTest(char *data, int2 *data_int2, unsigned int len);

int main()
{
    const int len = 16;
    int2 i2[16];                    // cuda 内置的 int2 类型
    char str[len] = { 82, 111, 118, 118,121, 42, 97, 121, 124, 118, 110, 56, 10, 10, 10,  10};    
    for (int i = 0; i < len; i++)
    {
        i2[i].x = str[i];
        i2[i].y = 10;
    }    
    runTest(str, i2, len);

    getchar();
    return 0;
}

● 输出结果：

ref char*: Hello World.
ref int2 :Hello World.
gpu char*: Hello World.
gpu int2 :Hello World.

▶ 涨姿势：

● cuda 内置的 int2 类型，整数有序对。涉及的定义如下：

#define __cuda_builtin_vector_align8(tag, members)  
    struct __device_builtin__ __align__(8) tag      
    {                                               
        members                                     
    }

__cuda_builtin_vector_align8(int2, int x; int y;);

typedef __device_builtin__ struct int2 int2;

● 警告函数和错误检查函数

#define assert(expression) (void)                                                                   
(                                                                                                   
(!!(expression)) || (_wassert(_CRT_WIDE(#expression), _CRT_WIDE(__FILE__), (unsigned)(__LINE__)), 0)
)

#define getLastCudaError(msg) __getLastCudaError (msg, __FILE__, __LINE__)

inline void __getLastCudaError(const char *errorMessage, const char *file, const int line)
{
    cudaError_t err = cudaGetLastError();
    if (cudaSuccess != err)
    {
        fprintf(stderr, "%s(%i) : getLastCudaError() CUDA error : %s : (%d) %s.
",
            file, line, errorMessage, (int)err, cudaGetErrorString(err));
        DEVICE_RESET
        exit(EXIT_FAILURE);
    }
}