0_Simple__inlinePTX + 0_Simple__inlinePTX_nvrtc

在核函数代码中加入并行线程执行（Parallel Thread eXecution，PTX），通过汇编指令获取得有关线程束的信息。并且在静态代码和运行时编译两种条件下使用。

▶ 源代码：静态使用

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_functions.h>
#include <helper_cuda.h>

__global__ void sequence_gpu(int *d_ptr, int length)
{
    int elemID = blockIdx.x * blockDim.x + threadIdx.x;

    if (elemID < length)
    {
        unsigned int laneid;
        asm("mov.u32 %0, %%laneid;" : "=r"(laneid));// 获取当前线程在线程束中的编号
        d_ptr[elemID] = laneid;
    }
}

void sequence_cpu(int *h_ptr, int length)
{
    for (int elemID=0; elemID<length; elemID++)
        h_ptr[elemID] = elemID % 32;
}

int main(int argc, char **argv)
{
    printf("CUDA inline PTX assembler sample
");

    const int N = 1000;

    int dev = findCudaDevice(argc, (const char **) argv);
    if (dev == -1)
        return EXIT_FAILURE;

    int *d_ptr;
    cudaMalloc(&d_ptr, N * sizeof(int));
    int *h_ptr;
    cudaMallocHost(&h_ptr, N * sizeof(int));

    dim3 cudaBlockSize(256,1,1);
    dim3 cudaGridSize((N + cudaBlockSize.x - 1) / cudaBlockSize.x, 1, 1);
    sequence_gpu<<<cudaGridSize, cudaBlockSize>>>(d_ptr, N);
    cudaGetLastError();
    cudaDeviceSynchronize();

    sequence_cpu(h_ptr, N);

    int *h_d_ptr;
    cudaMallocHost(&h_d_ptr, N *sizeof(int));
    cudaMemcpy(h_d_ptr, d_ptr, N *sizeof(int), cudaMemcpyDeviceToHost);

    bool bValid = true;

    for (int i=0; i<N && bValid; i++)
    {
        if (h_ptr[i] != h_d_ptr[i])
            bValid = false;
    }

    printf("Test %s.
", bValid ? "Successful" : "Failed");

    cudaFree(d_ptr);
    cudaFreeHost(h_ptr);
    cudaFreeHost(h_d_ptr);

    getchar();
    return bValid ? EXIT_SUCCESS: EXIT_FAILURE;
}

▶ 源代码：运行时编译

/*inlinePTX_kernel.cu*/
extern "C" __global__ void sequence_gpu(int *d_ptr, int length) 
{ 
    int elemID = blockIdx.x * blockDim.x + threadIdx.x; 
    if (elemID < length)
    {
        unsigned int laneid;
        asm("mov.u32 %0, %%laneid;" : "=r"(laneid));
        d_ptr[elemID] = laneid;
    }
}

/*inlinePTX.cpp*/
#include <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <nvrtc_helper.h>
#include <helper_functions.h>

void sequence_cpu(int *h_ptr, int length)
{
    for (int elemID=0; elemID<length; elemID++)
        h_ptr[elemID] = elemID % 32;
}

int main(int argc, char **argv)
{
    printf("CUDA inline PTX assembler sample
");

    char *ptx, *kernel_file;
    size_t ptxSize;

    kernel_file = sdkFindFilePath("inlinePTX_kernel.cu", argv[0]);
    compileFileToPTX(kernel_file, 0, NULL, &ptx, &ptxSize);
    CUmodule module = loadPTX(ptx, argc, argv);
    CUfunction kernel_addr;
    cuModuleGetFunction(&kernel_addr, module, "sequence_gpu");

    const int N = 1000;
    int *h_ptr = (int *)malloc(N * sizeof(int));

    dim3 cudaBlockSize(256,1,1);
    dim3 cudaGridSize((N + cudaBlockSize.x - 1) / cudaBlockSize.x, 1, 1);
    CUdeviceptr d_ptr;
    cuMemAlloc(&d_ptr, N * sizeof(int));

    void *arr[] = { (void *)&d_ptr, (void *)&N };
    cuLaunchKernel(kernel_addr,
                   cudaGridSize.x, cudaGridSize.y, cudaGridSize.z, 
                   cudaBlockSize.x, cudaBlockSize.y, cudaBlockSize.z,
                   0, 0, &arr[0], 0);

    cuCtxSynchronize();
    sequence_cpu(h_ptr, N);
    int *h_d_ptr = (int *)malloc(N * sizeof(int));;
    cuMemcpyDtoH(h_d_ptr, d_ptr, N *sizeof(int));

    bool bValid = true;
    for (int i=0; i<N && bValid; i++)
    {
        if (h_ptr[i] != h_d_ptr[i])
            bValid = false;
    }

    printf("Test %s.
", bValid ? "Successful" : "Failed");
    cuMemFree(d_ptr);

    getchar();
    return bValid ? EXIT_SUCCESS: EXIT_FAILURE;
}

▶ 输出结果：

CUDA inline PTX assembler sample
GPU Device 0: "GeForce GTX 1070" with compute capability 6.1

Test Successful.

▶ 涨姿势：

● 获取当前线程在线程束中的编号，即同意先乘数中的线程分别获得值 0 ~ 31

asm("mov.u32 %0, %%laneid;" : "=r"(laneid));