CbC/CbC_llvm: docs/NVPTXUsage.rst annotate

annotate docs/NVPTXUsage.rst @ 18:5abb684876c4

report error when code segment's tail call flag was changed false

author	Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp>
date	Sun, 06 Oct 2013 00:17:31 +0900
parents	9ad51c7bc036
children

rev	line source
3 9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	1 =============================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	2 User Guide for NVPTX Back-end
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	3 =============================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	4
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	5 .. contents::
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	6 :local:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	7 :depth: 3
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	8
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	9
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	10 Introduction
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	11 ============
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	12
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	13 To support GPU programming, the NVPTX back-end supports a subset of LLVM IR
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	14 along with a defined set of conventions used to represent GPU programming
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	15 concepts. This document provides an overview of the general usage of the back-
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	16 end, including a description of the conventions used and the set of accepted
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	17 LLVM IR.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	18
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	19 .. note::
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	20
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	21 This document assumes a basic familiarity with CUDA and the PTX
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	22 assembly language. Information about the CUDA Driver API and the PTX assembly
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	23 language can be found in the `CUDA documentation
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	24 <http://docs.nvidia.com/cuda/index.html>`_.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	25
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	26
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	27
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	28 Conventions
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	29 ===========
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	30
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	31 Marking Functions as Kernels
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	32 ----------------------------
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	33
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	34 In PTX, there are two types of functions: device functions, which are only
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	35 callable by device code, and kernel functions, which are callable by host
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	36 code. By default, the back-end will emit device functions. Metadata is used to
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	37 declare a function as a kernel function. This metadata is attached to the
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	38 ``nvvm.annotations`` named metadata object, and has the following format:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	39
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	40 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	41
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	42 !0 = metadata !{<function-ref>, metadata !"kernel", i32 1}
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	43
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	44 The first parameter is a reference to the kernel function. The following
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	45 example shows a kernel function calling a device function in LLVM IR. The
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	46 function ``@my_kernel`` is callable from host code, but ``@my_fmad`` is not.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	47
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	48 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	49
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	50 define float @my_fmad(float %x, float %y, float %z) {
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	51 %mul = fmul float %x, %y
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	52 %add = fadd float %mul, %z
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	53 ret float %add
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	54 }
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	55
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	56 define void @my_kernel(float* %ptr) {
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	57 %val = load float* %ptr
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	58 %ret = call float @my_fmad(float %val, float %val, float %val)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	59 store float %ret, float* %ptr
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	60 ret void
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	61 }
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	62
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	63 !nvvm.annotations = !{!1}
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	64 !1 = metadata !{void (float) @my_kernel, metadata !"kernel", i32 1}
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	65
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	66 When compiled, the PTX kernel functions are callable by host-side code.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	67
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	68
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	69 Address Spaces
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	70 --------------
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	71
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	72 The NVPTX back-end uses the following address space mapping:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	73
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	74 ============= ======================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	75 Address Space Memory Space
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	76 ============= ======================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	77 0 Generic
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	78 1 Global
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	79 2 Internal Use
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	80 3 Shared
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	81 4 Constant
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	82 5 Local
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	83 ============= ======================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	84
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	85 Every global variable and pointer type is assigned to one of these address
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	86 spaces, with 0 being the default address space. Intrinsics are provided which
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	87 can be used to convert pointers between the generic and non-generic address
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	88 spaces.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	89
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	90 As an example, the following IR will define an array ``@g`` that resides in
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	91 global device memory.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	92
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	93 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	94
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	95 @g = internal addrspace(1) global [4 x i32] [ i32 0, i32 1, i32 2, i32 3 ]
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	96
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	97 LLVM IR functions can read and write to this array, and host-side code can
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	98 copy data to it by name with the CUDA Driver API.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	99
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	100 Note that since address space 0 is the generic space, it is illegal to have
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	101 global variables in address space 0. Address space 0 is the default address
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	102 space in LLVM, so the ``addrspace(N)`` annotation is required for global
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	103 variables.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	104
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	105
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	106 NVPTX Intrinsics
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	107 ================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	108
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	109 Address Space Conversion
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	110 ------------------------
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	111
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	112 '``llvm.nvvm.ptr.*.to.gen``' Intrinsics
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	113 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	114
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	115 Syntax:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	116 """""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	117
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	118 These are overloaded intrinsics. You can use these on any pointer types.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	119
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	120 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	121
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	122 declare i8* @llvm.nvvm.ptr.global.to.gen.p0i8.p1i8(i8 addrspace(1)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	123 declare i8* @llvm.nvvm.ptr.shared.to.gen.p0i8.p3i8(i8 addrspace(3)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	124 declare i8* @llvm.nvvm.ptr.constant.to.gen.p0i8.p4i8(i8 addrspace(4)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	125 declare i8* @llvm.nvvm.ptr.local.to.gen.p0i8.p5i8(i8 addrspace(5)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	126
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	127 Overview:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	128 """""""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	129
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	130 The '``llvm.nvvm.ptr.*.to.gen``' intrinsics convert a pointer in a non-generic
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	131 address space to a generic address space pointer.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	132
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	133 Semantics:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	134 """"""""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	135
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	136 These intrinsics modify the pointer value to be a valid generic address space
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	137 pointer.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	138
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	139
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	140 '``llvm.nvvm.ptr.gen.to.*``' Intrinsics
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	141 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	142
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	143 Syntax:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	144 """""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	145
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	146 These are overloaded intrinsics. You can use these on any pointer types.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	147
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	148 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	149
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	150 declare i8* @llvm.nvvm.ptr.gen.to.global.p1i8.p0i8(i8 addrspace(1)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	151 declare i8* @llvm.nvvm.ptr.gen.to.shared.p3i8.p0i8(i8 addrspace(3)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	152 declare i8* @llvm.nvvm.ptr.gen.to.constant.p4i8.p0i8(i8 addrspace(4)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	153 declare i8* @llvm.nvvm.ptr.gen.to.local.p5i8.p0i8(i8 addrspace(5)*)
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	154
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	155 Overview:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	156 """""""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	157
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	158 The '``llvm.nvvm.ptr.gen.to.*``' intrinsics convert a pointer in the generic
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	159 address space to a pointer in the target address space. Note that these
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	160 intrinsics are only useful if the address space of the target address space of
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	161 the pointer is known. It is not legal to use address space conversion
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	162 intrinsics to convert a pointer from one non-generic address space to another
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	163 non-generic address space.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	164
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	165 Semantics:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	166 """"""""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	167
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	168 These intrinsics modify the pointer value to be a valid pointer in the target
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	169 non-generic address space.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	170
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	171
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	172 Reading PTX Special Registers
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	173 -----------------------------
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	174
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	175 '``llvm.nvvm.read.ptx.sreg.*``'
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	176 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	177
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	178 Syntax:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	179 """""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	180
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	181 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	182
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	183 declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	184 declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	185 declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	186 declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	187 declare i32 @llvm.nvvm.read.ptx.sreg.ntid.y()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	188 declare i32 @llvm.nvvm.read.ptx.sreg.ntid.z()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	189 declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	190 declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	191 declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	192 declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	193 declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	194 declare i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	195 declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	196
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	197 Overview:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	198 """""""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	199
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	200 The '``@llvm.nvvm.read.ptx.sreg.*``' intrinsics provide access to the PTX
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	201 special registers, in particular the kernel launch bounds. These registers
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	202 map in the following way to CUDA builtins:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	203
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	204 ============ =====================================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	205 CUDA Builtin PTX Special Register Intrinsic
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	206 ============ =====================================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	207 ``threadId`` ``@llvm.nvvm.read.ptx.sreg.tid.*``
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	208 ``blockIdx`` ``@llvm.nvvm.read.ptx.sreg.ctaid.*``
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	209 ``blockDim`` ``@llvm.nvvm.read.ptx.sreg.ntid.*``
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	210 ``gridDim`` ``@llvm.nvvm.read.ptx.sreg.nctaid.*``
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	211 ============ =====================================
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	212
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	213
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	214 Barriers
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	215 --------
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	216
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	217 '``llvm.nvvm.barrier0``'
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	218 ^^^^^^^^^^^^^^^^^^^^^^^^^^^
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	219
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	220 Syntax:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	221 """""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	222
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	223 .. code-block:: llvm
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	224
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	225 declare void @llvm.nvvm.barrier0()
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	226
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	227 Overview:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	228 """""""""
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	229
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	230 The '``@llvm.nvvm.barrier0()``' intrinsic emits a PTX ``bar.sync 0``
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	231 instruction, equivalent to the ``__syncthreads()`` call in CUDA.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	232
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	233
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	234 Other Intrinsics
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	235 ----------------
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	236
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	237 For the full set of NVPTX intrinsics, please see the
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	238 ``include/llvm/IR/IntrinsicsNVVM.td`` file in the LLVM source tree.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	239
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	240
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	241 Executing PTX
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	242 =============
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	243
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	244 The most common way to execute PTX assembly on a GPU device is to use the CUDA
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	245 Driver API. This API is a low-level interface to the GPU driver and allows for
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	246 JIT compilation of PTX code to native GPU machine code.
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	247
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	248 Initializing the Driver API:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	249
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	250 .. code-block:: c++
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	251
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	252 CUdevice device;
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	253 CUcontext context;
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	254
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	255 // Initialize the driver API
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	256 cuInit(0);
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	257 // Get a handle to the first compute device
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	258 cuDeviceGet(&device, 0);
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	259 // Create a compute device context
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	260 cuCtxCreate(&context, 0, device);
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	261
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	262 JIT compiling a PTX string to a device binary:
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	263
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	264 .. code-block:: c++
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	265
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	266 CUmodule module;
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	267 CUfunction funcion;
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	268
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	269 // JIT compile a null-terminated PTX string
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	270 cuModuleLoadData(&module, (void*)PTXString);
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	271
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	272 // Get a handle to the "myfunction" kernel function
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	273 cuModuleGetFunction(&function, module, "myfunction");
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	274
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	275 For full examples of executing PTX assembly, please see the `CUDA Samples
9ad51c7bc036 1st commit. remove git dir and add all files. Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	276 <https://developer.nvidia.com/cuda-downloads>`_ distribution.

Mercurial > hg > CbC > CbC_llvm

annotate docs/NVPTXUsage.rst @ 18:5abb684876c4