User Guide for NVPTX Back-end — LLVM 3.3 documentation

# User Guide for NVPTX Back-end

- Introduction
- Conventions
  - Marking Functions as Kernels
  - Address Spaces
- NVPTX Intrinsics
  - Address Space Conversion
    - ‘`llvm.nvvm.ptr.*.to.gen`‘ Intrinsics
    - ‘`llvm.nvvm.ptr.gen.to.*`‘ Intrinsics
  - Reading PTX Special Registers
    - ‘`llvm.nvvm.read.ptx.sreg.*`‘
  - Barriers
    - ‘`llvm.nvvm.barrier0`‘
  - Other Intrinsics
- Executing PTX

## Introduction

To support GPU programming, the NVPTX back-end supports a subset of LLVM IR along with a defined set of conventions used to represent GPU programming concepts. This document provides an overview of the general usage of the back- end, including a description of the conventions used and the set of accepted LLVM IR.

Note

This document assumes a basic familiarity with CUDA and the PTX assembly language. Information about the CUDA Driver API and the PTX assembly language can be found in the CUDA documentation.

## Conventions

### Marking Functions as Kernels

In PTX, there are two types of functions: device functions, which are only callable by device code, and kernel functions, which are callable by host code. By default, the back-end will emit device functions. Metadata is used to declare a function as a kernel function. This metadata is attached to the `nvvm.annotations` named metadata object, and has the following format:

    !0 = metadata !{<function-ref>, metadata !"kernel", i32 1}

The first parameter is a reference to the kernel function. The following example shows a kernel function calling a device function in LLVM IR. The function `@my_kernel` is callable from host code, but `@my_fmad` is not.

    define float @my_fmad(float %x, float %y, float %z) {
      %mul = fmul float %x, %y
      %add = fadd float %mul, %z
      ret float %add
    }

    define void @my_kernel(float* %ptr) {
      %val = load float* %ptr
      %ret = call float @my_fmad(float %val, float %val, float %val)
      store float %ret, float* %ptr
      ret void
    }

    !nvvm.annotations = !{!1}
    !1 = metadata !{void (float*)* @my_kernel, metadata !"kernel", i32 1}

When compiled, the PTX kernel functions are callable by host-side code.

### Address Spaces

The NVPTX back-end uses the following address space mapping:

> Address Space | Memory Space
> --- | ---
> 0 | Generic
> 1 | Global
> 2 | Internal Use
> 3 | Shared
> 4 | Constant
> 5 | Local

Every global variable and pointer type is assigned to one of these address spaces, with 0 being the default address space. Intrinsics are provided which can be used to convert pointers between the generic and non-generic address spaces.

As an example, the following IR will define an array `@g` that resides in global device memory.

    @g = internal addrspace(1) global [4 x i32] [ i32 0, i32 1, i32 2, i32 3 ]

LLVM IR functions can read and write to this array, and host-side code can copy data to it by name with the CUDA Driver API.

Note that since address space 0 is the generic space, it is illegal to have global variables in address space 0. Address space 0 is the default address space in LLVM, so the `addrspace(N)` annotation is required for global variables.

## NVPTX Intrinsics

### Address Space Conversion

#### ‘`llvm.nvvm.ptr.*.to.gen`‘ Intrinsics

##### Syntax:

These are overloaded intrinsics. You can use these on any pointer types.

    declare i8* @llvm.nvvm.ptr.global.to.gen.p0i8.p1i8(i8 addrspace(1)*)
    declare i8* @llvm.nvvm.ptr.shared.to.gen.p0i8.p3i8(i8 addrspace(3)*)
    declare i8* @llvm.nvvm.ptr.constant.to.gen.p0i8.p4i8(i8 addrspace(4)*)
    declare i8* @llvm.nvvm.ptr.local.to.gen.p0i8.p5i8(i8 addrspace(5)*)

##### Overview:

The ‘`llvm.nvvm.ptr.*.to.gen`‘ intrinsics convert a pointer in a non-generic address space to a generic address space pointer.

##### Semantics:

These intrinsics modify the pointer value to be a valid generic address space pointer.

#### ‘`llvm.nvvm.ptr.gen.to.*`‘ Intrinsics

##### Syntax:

These are overloaded intrinsics. You can use these on any pointer types.

    declare i8* @llvm.nvvm.ptr.gen.to.global.p1i8.p0i8(i8 addrspace(1)*)
    declare i8* @llvm.nvvm.ptr.gen.to.shared.p3i8.p0i8(i8 addrspace(3)*)
    declare i8* @llvm.nvvm.ptr.gen.to.constant.p4i8.p0i8(i8 addrspace(4)*)
    declare i8* @llvm.nvvm.ptr.gen.to.local.p5i8.p0i8(i8 addrspace(5)*)

##### Overview:

The ‘`llvm.nvvm.ptr.gen.to.*`‘ intrinsics convert a pointer in the generic address space to a pointer in the target address space. Note that these intrinsics are only useful if the address space of the target address space of the pointer is known. It is not legal to use address space conversion intrinsics to convert a pointer from one non-generic address space to another non-generic address space.

##### Semantics:

These intrinsics modify the pointer value to be a valid pointer in the target non-generic address space.

### Reading PTX Special Registers

#### ‘`llvm.nvvm.read.ptx.sreg.*`‘

##### Syntax:

    declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
    declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
    declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
    declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
    declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
    declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
    declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
    declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
    declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
    declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
    declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
    declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
    declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()

##### Overview:

The ‘`@llvm.nvvm.read.ptx.sreg.*`‘ intrinsics provide access to the PTX special registers, in particular the kernel launch bounds. These registers map in the following way to CUDA builtins:

> CUDA Builtin | PTX Special Register Intrinsic
> --- | ---
> `threadId` | `@llvm.nvvm.read.ptx.sreg.tid.*`
> `blockIdx` | `@llvm.nvvm.read.ptx.sreg.ctaid.*`
> `blockDim` | `@llvm.nvvm.read.ptx.sreg.ntid.*`
> `gridDim` | `@llvm.nvvm.read.ptx.sreg.nctaid.*`

### Barriers

#### ‘`llvm.nvvm.barrier0`‘

##### Syntax:

    declare void @llvm.nvvm.barrier0()

##### Overview:

The ‘`@llvm.nvvm.barrier0()`‘ intrinsic emits a PTX `bar.sync 0` instruction, equivalent to the `__syncthreads()` call in CUDA.

### Other Intrinsics

For the full set of NVPTX intrinsics, please see the `include/llvm/IR/IntrinsicsNVVM.td` file in the LLVM source tree.

## Executing PTX

The most common way to execute PTX assembly on a GPU device is to use the CUDA Driver API. This API is a low-level interface to the GPU driver and allows for JIT compilation of PTX code to native GPU machine code.

Initializing the Driver API:

    CUdevice device;
    CUcontext context;

    // Initialize the driver API
    cuInit(0);
    // Get a handle to the first compute device
    cuDeviceGet(&device, 0);
    // Create a compute device context
    cuCtxCreate(&context, 0, device);

JIT compiling a PTX string to a device binary:

    CUmodule module;
    CUfunction funcion;

    // JIT compile a null-terminated PTX string
    cuModuleLoadData(&module, (void*)PTXString);

    // Get a handle to the "myfunction" kernel function
    cuModuleGetFunction(&function, module, "myfunction");

For full examples of executing PTX assembly, please see the CUDA Samples distribution.