C++for OpenCL 1.0和2021编程语言文档

The C++ for OpenCL 1.0 and 2021 Programming Language Documentation

C++for OpenCL 1.0和2021编程语言文档

Khronos® OpenCL Working Group

Version DocRev2021.12, Mon, 17 Apr 2023 17:53:43 +0000: from git branch: commit: 8a5ddd4101834b5e376d286fd1a8c02d5c235e7a tag: v3.0.14

Copyright 2019-2023 The Khronos Group.

版权所有2019-2023Khronos 集团。

Khronos licenses this file to you under the Creative Commons Attribution 4.0 International (CC BY 4.0) License (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at https://creativecommons.org/licenses/by/4.0/

Khronos根据知识共享署名4.0国际（CC BY 4.0）许可证（以下简称“许可证”）将本文件授权给您；除非符合许可证的规定，否则不得使用此文件。可以在以下网址获取许可证副本：https://creativecommons.org/licenses/by/4.0/

Unless required by applicable law or agreed to in writing, material distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. If all or a portion of this material is re-used, notice substantially similar to the following must be included:

除非适用法律要求或书面同意，否则根据许可证分发的材料将按“原样”分发，不存在任何明示或暗示的担保或条件。如果全部或部分材料被重复使用，则必须包括与以下内容基本相似的通知：

This documentation includes material developed at The Khronos Group (http://www.khronos.org/). Khronos supplied such material on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied, under the terms of the Creative Commons Attribution 4.0 International (CC BY 4.0) License (the "License"), available at https://creativecommons.org/licenses/ by/4.0/. All use of such material is governed by the term of the License. Khronos bears no responsibility whatsoever for additions or modifications to its material.

本文件包括Khronos集团开发的材料(http://www.khronos.org/)。Khronos根据知识共享署名4.0国际（CC BY 4.0）许可证（“许可证”）的条款“按原样”提供此类材料，不存在任何明示或暗示的担保或条件，可在https://creativecommons.org/licenses/通过/4.0/。此类材料的所有使用均受许可条款的约束。Khronos对其材料的添加或修改不承担任何责任。

Khronos is a registered trademark, and OpenCL is a trademark of Apple Inc. and used under license by Khronos. All other product names, trademarks, and/or company names are used solely for identification and belong to their respective owners.

Khronos是一个注册商标，OpenCL是苹果股份有限公司的商标，由Khronos授权使用。所有其他产品名称、商标和/或公司名称仅用于识别，属于其各自所有者。

Chapter 1. Introduction

第1章 介绍

This language is built on top of OpenCL C 3.0 unified and C++17 enabling most of regular C++ features in OpenCL kernel code. Most functionality from C++ and OpenCL C is inherited. Since both OpenCL C and C++ are derived from C and moreover C++ is almost fully backward compatible with C, the main design principle of C++ for OpenCL is to reapply existing OpenCL concepts to C++. Therefore, it is important to refer to The OpenCL Specification, Version 3.0 section 3.2 and section 3.3 detailing fundamental differences of OpenCL execution and memory models from the conventional C and C++ view.

该语言建立在OpenCL C 3.0统一版和C++17之上，支持OpenCL内核代码中的大多数常规C++功能。C++和OpenCLC的大部分功能都是继承的。由于OpenCL C和C++都是从C派生而来的，而且C++几乎与C完全向后兼容，因此用于OpenCL的C++的主要设计原则是将现有的OpenCL概念重新应用于C++。因此，重要的是参考OpenCL规范3.0版第3.2节和第3.3节，详细说明了OpenCL执行和内存模型与传统C和C++视图的基本差异。

This document describes the programming language in detail. It is not structured as a standalone document, but rather as an addition to OpenCL C 3.0 unified specification defined in The OpenCL C Specification, Version 3.0 and C++17 defined in The C++17 Specification. Where necessary this document refers to the specifications of those languages accordingly. A full understanding of C++ for OpenCL requires familiarity with the specifications or other documentation of both languages that C++ for OpenCL is built upon.

本文档详细介绍了编程语言。它不是作为一个独立的文档构建的，而是对OpenCL C规范中定义的OpenCL C 3.0统一规范的补充，版本3.0和C++17规范中定义了C++17。如有必要，本文件相应地参考这些语言的规范。对C++for OpenCL的全面理解需要熟悉C++for OpenCL所基于的两种语言的规范或其他文档。

The description of C++ for OpenCL starts from highlighting the differences to OpenCL C and the differences to C++.

对OpenCL的C++的描述从强调与OpenCL C的差异和与C++的差异开始。

The majority of content covers the behavior that is not documented in the OpenCL C 3.0 s6 and C++17 specifications. This is mainly related to interactions between OpenCL and C++ language features.

大部分内容涵盖了OpenCL C 3.0 s6和C++17规范中未记录的行为。这主要与OpenCL和C++语言特性之间的交互有关。

This document describes C++ for OpenCL language

本文档描述了用于OpenCL语言的C++

• version 1.0 that is backward compatible with OpenCL 2.0; and

•与OpenCL 2.0向后兼容的1.0版本；和

• version 2021 that is backward compatible with OpenCL 3.0.

•与OpenCL 3.0向后兼容的2021版。

C++ for OpenCL 2021 is described in this document as a provisional language version. While no large changes are envisioned in the future, some minor aspects might not remain identical in its final release

用于OpenCL 2021的C++在本文档中被描述为临时语言版本。虽然未来不会有大的变化，但在最终版本中，一些小方面可能不会完全相同

 Chapter 2. Version differences

第2章 版本差异

The main difference between C++ for OpenCL version 1.0 and version 2021 comes from the difference between OpenCL 2.0 and OpenCL 3.0 with which they are respectively compatible. Support for some features of C++ for OpenCL 1.0 has become optional as described in OpenCL 3.0 s6.2.1. Predefined feature macros from OpenCL C 3.0 can be used to detect which optional features are present.

用于OpenCL 1.0版和2021版的C++之间的主要差异来自于它们分别兼容的OpenCL 2.0和OpenCL 3.0之间的差异。如OpenCL 3.0 s6.2.1中所述，对OpenCL 1.0的C++某些功能的支持已成为可选的。OpenCL C 3.0中预定义的功能宏可用于检测存在哪些可选功能。

This impacts some C++ specific semantics mainly due to optionality of the generic address space (i.e. __opencl_c_generic_address_space feature) or program scope variables (i.e. __opencl_c_program_scope_global_variables feature).

这影响了一些C++特定的语义，主要是由于通用地址空间（即__opencl_C_generic_address_space特性）或程序范围变量（即__open C_program_scope_global_variables特性）的可选性。

Chapter 3. The C++ for OpenCL Programming Language

第3章 面向OpenCL程序设计语言的C++

This programming language inherits features from OpenCL C 3.0 s6 as well as C++17. Detailed aspects of OpenCL and C++ are not described in this document as they can be found in their official specifications.

这种编程语言继承了OpenCL C 3.0 s6以及C++17的特性。OpenCL和C++的详细方面在本文档中没有描述，因为它们可以在其官方规范中找到。

This section documents various language features of C++ for OpenCL that are not covered in neither OpenCL nor C++ specifications, in particular:

本节记录了OpenCL的C++的各种语言功能，这些功能在OpenCL和C++规范中都没有涵盖，特别是：

• any behavior that deviates from C++17;

•任何偏离C++17的行为；

• any behavior that deviates from OpenCL C 2.0 or 3.0;

•任何偏离OpenCL C 2.0或3.0的行为；

• any behavior that is not governed by OpenCL C and C++.

•任何不受OpenCL C和C++控制的行为。

All language extensions to OpenCL C are applicable to C++ for OpenCL.

OpenCL C的所有语言扩展都适用于OpenCL的C++。

• Extensions to OpenCL C 2.0 or earlier versions apply to C++ for OpenCL version 1.0.

•对OpenCL C 2.0或更早版本的扩展适用于OpenCL 1.0版的C++。

• Extensions to OpenCL C 3.0 or earlier versions except for OpenCL C 2.0, apply to C++ for OpenCL 2021.

•除OpenCL C 2.0外，对OpenCL C 3.0或更早版本的扩展适用于OpenCL 2021的C++。

3.1. Difference to C++

3.1.与C的差异++

C++ for OpenCL supports the majority of standard C++17 features, however, there are some differences that are documented in this section.

用于OpenCL的C++支持大多数标准C++17功能，但是，本节中记录了一些差异。

3.1.1. Restrictions to C++ features

3.1.1.对C++功能的限制

The following C++ language features are not supported:

不支持以下C++语言功能：

• Virtual functions (C++17 [class.virtual]);

•虚拟函数（C++17[class.Virtual]）；

• References to functions (C++17 [class.mfct]);

•对函数的引用（C++17[class.mfct]）；

• Pointers to class member functions (in addition to the regular non-member functions that are already restricted in OpenCL C);

•指向类成员函数的指针（除了OpenCL C中已经限制的常规非成员函数之外）；

• Exceptions (C++17 [except]);

•例外情况（C++17[除外]）；

• dynamic_cast operator (C++17 [expr.dynamic.cast]);

•dynamic_cast运算符（C++17[expr.dynamic.cast]）；

• Non-placement new/delete operators (C++17 [expr.new]/[expr.delete]);

•非放new/delete运算符（C++17[expr.new]/[expr.delete]）；

• thread_local storage class specifier (C++17 [basic.stc.thread]);

•thread_local存储类说明符（C++17[basic.stc.thread]）；

• Standard C++ libraries (C++17 [library]).

•标准C++库（C++17[库]）。

Simultaneous initialization of static local objects performed by different work-items is not guaranteed to be free from race-conditions. Whether an implementation provides such a guarantee is indicated by the presence of the __cpp_threadsafe_static_init feature test macro[1] .

由不同工作项执行的静态本地对象的同时初始化不能保证不受竞争条件的影响。实现是否提供这样的保证由__cpp_threadsafe_static_init功能测试宏[1]的存在来指示。

The list above only contains extra restrictions that are not detailed in OpenCL C specification. As OpenCL restricts a number of C features, the same restrictions are inherited by C++ for OpenCL. The detailed list of C feature restrictions is provided in OpenCL C 3.0 s6.11.

上面的列表只包含OpenCL C规范中没有详细说明的额外限制。由于OpenCL限制了许多C特性，C++也继承了同样的限制。OpenCL C 3.0 s6.11中提供了C功能限制的详细列表。

3.2. Difference to OpenCL C

3.2.与OpenCL C的区别

C++ for OpenCL provides backwards compatibility with OpenCL C for the majority of features. However, there are a number of exceptions that are described in this section. Some of them come from the nature of C++ but others are due to improvements in OpenCL features. Most of such improvements do not invalidate code written in OpenCL C, but simply provide extra functionality.

用于OpenCL的C++为大多数功能提供了与OpenCL C的向后兼容性。但是，本节中介绍了一些例外情况。其中一些来自C++的本质，但另一些则是由于OpenCL功能的改进。大多数这样的改进不会使用OpenCLC编写的代码失效，只是提供了额外的功能。

3.2.1. C++ related differences

3.2.1.C++相关差异

C++ for OpenCL is a different language to OpenCL C and it is derived from C++ inheriting C++'s fundamental design principles. Hence C++ for OpenCL deviates from OpenCL C in the same areas where C++ deviates from C. This results in a more helpful language for developers and facilitates improvements in compilation tools without substantially increasing their complexity.

用于OpenCL的C++是一种不同于OpenCLC的语言，它源于继承C++基本设计原则的C++。因此，用于OpenCL的C++在C++偏离C的相同领域偏离了OpenCLC。这为开发人员带来了更有用的语言，并促进了编译工具的改进，而不会显著增加它们的复杂性。

3.2.1.1. Implicit conversions

3.2.1.1.隐式转换

C++ is much stricter about conversions between types, especially those that are performed implicitly by the compiler. For example it is not possible to convert a const object to non-const implicitly. For details please refer to C++17 [conv].

C++对类型之间的转换要严格得多，尤其是那些由编译器隐式执行的转换。例如，不可能隐式地将常量对象转换为非常量。有关详细信息，请参阅C++17[conv]。

void foo(){
  const int *ptrconst;
  int *ptr = ptrconst; // invalid initialization discards const qualifier.
                       // 无效的初始化将丢弃常量限定符。
}

The same applies to narrowing conversions in initialization lists (C++17 [dcl.init.list]).

这同样适用于缩小初始化列表（C++17[dcl.init.list]）中的转换范围。

struct mytype {
    int i;
};
void foo(uint par){
    mytype var = {
        .i = par // narrowing from uint to int is disallowed.
                 //不允许从uint缩小到int。
    };
}

Some compilers allow silencing this error using a flag (e.g. in Clang -Wno-error=c++11-narrowing can be used).

一些编译器允许使用标志来消除此错误（例如，在Clang-Wno error=c++11中，可以使用窄化）。

Among other common conversions that will not be compiled in C++ mode there are pointer to integer or integer to pointer type conversions.

在其他不会以C++模式编译的常见转换中，有指针到整数或整数到指针类型的转换。

void foo(){
  int *ptr;
  int i = ptr; // incompatible pointer to integer conversion.
               //不兼容的指针到整数转换。
}

3.2.1.2. Null pointer constant

3.2.1.2.空指针常量

In C and OpenCL C the null pointer constant is defined using other language features as it is not represented explicitly i.e. commonly it is defined as

在C和OpenCL C中，空指针常量是使用其他语言特性定义的，因为它没有明确表示，即通常定义为

#define NULL ((void*)0)

In C++17 there is an explicit builtin pointer literal nullptr that should be used instead (C++17 [lex.nullptr]).

在C++17中，有一个显式的内置指针文本nullptr，应该使用它（C++17[lex.nullptr]）。

NULL macro definition in C++ for OpenCL follows C++17 [support.types.nullptr] where it is an implementation defined macro and it is not guaranteed to be the same as in OpenCL C. Reusing the definition of NULL from OpenCL C does not guarantee that any code with NULL is legal in C++ for OpenCL even if it is legal in OpenCL C.

C++中用于OpenCL的NULL宏定义遵循C++17[support.types.nullptr]，其中它是一个实现定义的宏，不保证与OpenCL C中的相同。从OpenCL C重复使用NULL的定义并不保证任何带有NULL的代码在C++中对于OpenCL都是合法的，即使它在OpenCL C是合法的。

#define NULL ((void*)0)
void foo(){
    int *ptr = NULL; // invalid initialization of int* with void*.
                     // 使用void*初始化int*无效。
}

To improve code portability and compatibility, implementations are encouraged to define NULL as an alias to pointer literal nullptr.

为了提高代码的可移植性和兼容性，鼓励实现将NULL定义为指针文本nullptr的别名。

3.2.1.3. Use of restrict

3.2.1.3.限制的使用

C++17 does not support restrict and therefore C++ for OpenCL can not support it either. Some compilers might provide extensions with some functionality of restrict in C++, e.g. __restrict in Clang.

C++17不支持restrict，因此OpenCL的C++也不能支持它。一些编译器可能会在C++中提供一些带有restrict功能的扩展，例如Clang中的__restrict。

This feature only affects optimizations and the source code can be modified by removing it. As a workaround to avoid manual modifications, macro substitutions can be used to either remove the keyword during the preprocessing by defining restrict as an empty macro or mapping it to another similar compiler features, e.g. __restrict in Clang. This can be done in headers or using -D compilation flag.

此功能只影响优化，并且可以通过删除源代码来修改源代码。作为避免手动修改的解决方法，可以使用宏替换在预处理期间通过将restrict定义为空宏来删除关键字，或者将其映射到另一个类似的编译器功能，例如Clang中的__restrict。这可以在头中完成，也可以使用-D编译标志。

3.2.1.4. Limitations of goto statements

3.2.1.4.goto语句的限制

C++ is more restrictive with respect to entering the scope of variables than C. It is not possible to jump forward over a variable declaration statement apart from some exceptions detailed in C++17 [stmt.dcl].

与C相比，C++在输入变量范围方面更具限制性。除了C++17[stmt.dcl]中详细介绍的一些例外情况外，无法跳过变量声明语句。

      if (cond)
          goto label;
      int n = foo();
label: // invalid: jumping forward over declaration of n.
       //无效：跳过n的声明。
  // ...

3.2.1.5. Ternary selection operator

3.2.1.5.三元选择算子

The ternary selection operator (?:) inherits its behaviour from both C++ and OpenCL C. It operates on three expressions (exp1 ? exp2 : exp3). If all three expressions are scalar values, the C++17 rules for ternary operator are followed. If the result is a vector value, then this is equivalent to calling select(exp3, exp2, exp1) as described in OpenCL C 3.0 s6.15.6. The rules from OpenCL C impose limitation that exp1 cannot be a vector of float values. However, exp1 can be evaluated to a scalar float as it is contextually convertible to bool in C++.

三元选择运算符（？：）继承了C++和OpenCL C的行为。它对三个表达式（exp1？exp2:exp3）进行操作。如果所有三个表达式都是标量值，则遵循三元运算符的C++17规则。如果结果是一个向量值，那么这相当于调用select（exp3，exp2，exp1），如OpenCL C 3.0 s615.6中所述。OpenCL C中的规则限制了exp1不能是浮点值的向量。然而，exp1可以被计算为标量浮点，因为它在C++中可以在上下文中转换为bool。

3.2.2. OpenCL specific difference

3.2.2 OpenCL特定差异

This section describes where C++ for OpenCL differs from OpenCL C in OpenCL specific behavior

本节描述OpenCL的C++与OpenCL C在OpenCL特定行为方面的不同之处

3.2.2.1. Variadic macros

3.2.2.1可变宏

C++ for OpenCL eliminates the restriction on variadic macros from OpenCL C 3.0 s6.11.f. Variadic macros can be used normally as per C++17 [cpp.replace].

用于OpenCL的C++消除了OpenCL C 3.0 s6.11.f对可变宏的限制。可变宏可以按照C++17[cpp.replace]正常使用。

3.2.2.2. Predefined macros

3.2.2.2.预定义宏

The macro __OPENCL_C_VERSION__ described in OpenCL C 3.0 s6.12, is not defined.

未定义OPENCL C 3.0 s6.12中描述的宏__OPENCL_C_VERSION__。

The following new predefined macros are added in C++ for OpenCL:

在C++for OpenCL中添加了以下新的预定义宏：

• __OPENCL_CPP_VERSION__ set to integer value reflecting the C++ for OpenCL version the translation unit is compiled for. The value 100 corresponds to the language version 1.0 and 202100 corresponds to the version 2021.

•__OPENCL_CPP_VERSION__设置为整数值，反映编译翻译单元的OPENCL版本的C++。值100对应于语言版本1.0，202100对应于版本2021。

• __CL_CPP_VERSION_1_0__ set to 100 and can be used for convenience instead of a literal.

•__CL_CPP_VERSION_1_0__设置为100，可用于方便而非文字。

• __CL_CPP_VERSION_2021__ set to 202100 and can be used for convenience instead of a literal.

•__CL_CPP_VERSION_2021__设置为202100，可用于方便而非文字。

3.2.2.3. Atomic operations

3.2.2.3.原子操作

C++ for OpenCL relaxes restriction from OpenCL C 3.0 s6.15.12 to atomic types allowing them to be used by builtin operators, and not only by builtin functions. This relaxation does not apply to C++ for OpenCL version 2021 if the sequential consistency memory model (i.e. __opencl_c_atomic_order_seq_cst feature) is not supported.

C++for OpenCL放宽了从OpenCL C 3.0 s615.12到原子类型的限制，允许它们由内置运算符使用，而不仅仅是由内置函数使用。如果顺序一致性内存模型（即__OpenCL_C_atomic_order_seq_cst功能）不受支持，则这种放松不适用于2021版OpenCL的C++。

Operators on atomic types behave as described in C++17 sections [atomics.types.int] [atomics.types.pointer] [atomics.types.float].

原子类型上的运算符的行为如C++17节[atomics.types.int][atomics.types.pointer][atomics.types.float]中所述。

// Assumes support of sequential consistency memory model.
//假定支持顺序一致性内存模型。
atomic_int acnt;
acnt++; // equivalent to atomic_fetch_add(&acnt, 1);
        //相当于atomic_fetch_add（&acnt，1）;

3.2.2.4. Use of Blocks

3.2.2.4.块的使用

Blocks that are defined in OpenCL C 3.0 s6.14 are not supported and their use might be replaced by lambdas (C++17 [expr.prim.lambda]) in future versions.

OpenCL C 3.0 s6.14中定义的块不受支持，在未来的版本中，它们的使用可能会被lambdas（C++17[expr.prim.lambda]）所取代。

The above implies that builtin functions using blocks, such as enqueue_kernel, are not supported in C++ for OpenCL.

上面的内容表明，在C++for OpenCL中不支持使用块的内置函数，例如enqueue_kernel。

3.3. Address spaces

3.3.地址空间

C++ for OpenCL inherits address space behavior from OpenCL C 3.0 s6.7.

用于OpenCL的C++继承了来自OpenCL C 3.0 s6.7的地址空间行为。

This section only documents behavior related to C++ features. For example, conversion rules are extended from the qualification conversion in C++17 [conv.qual] but the compatibility is determined using notation of sets and overlapping of address spaces from section 5.1.3 of The Embedded C Specification. For OpenCL kernel languages there are two main semantics depending on whether generic address space (OpenCL C 3.0 s6.7.5) is supported on not. The generic address space is always supported for

本节仅记录与C++功能相关的行为。例如，转换规则是从C++17[conv.qual]中的限定转换扩展而来的，但兼容性是使用嵌入式C规范第5.1.3节中的集合表示法和地址空间重叠来确定的。对于OpenCL内核语言，根据是否支持通用地址空间（OpenCL C 3.0 s6.7.5），有两种主要语义。始终支持的通用地址空间

• C++ for OpenCL 1.0;

•用于OpenCL 1.0的C++；

• C++ for OpenCL 2021 with the presence of __opencl_c_generic_address_space feature as explained in OpenCL C 3.0 s6.2.1.

•用于OpenCL 2021的C++，具有__OpenCL_C_generic_address_space功能，如OpenCL C 3.0 s6.2.1中所述。

If generic address space is not supported, qualification conversions with pointer type where address spaces differ are not allowed. If generic address space is supported, implicit conversions are allowed from a named address space (except for __constant) to generic address space. The reverse conversion is only allowed explicitly. The __constant address space does not overlap with any other, therefore, no valid conversion between __constant and any other address space exists. This is aligned with rules from OpenCL C 3.0 s6.7.9 and this logic regulates semantics described in this section.

如果不支持通用地址空间，则不允许使用地址空间不同的指针类型进行资格转换。如果支持通用地址空间，则允许从命名地址空间（__constant除外）到通用地址空间的隐式转换。仅明确允许反向转换。__constant地址空间与其他地址空间不重叠，因此，__constant和其他地址空间之间不存在有效的转换。这与OpenCL C 3.0 s6.7.9中的规则一致，并且该逻辑规范了本节中描述的语义。

3.3.1. Casts

3.3.1类型转换

C-style casts follow rules of OpenCL C 3.0 s6.7.9. Conversions of references and pointers to the generic address space can be done by any C++ cast operator (as an implicit conversion); converting from generic to named address space can only be done using the dedicated addrspace_cast operator. The addrspace_cast operator can only convert between address spaces for pointers and references and no other conversions are allowed to occur. Note that conversions between __constant and any other other address space are disallowed.

C样式强制转换遵循OpenCL C 3.0 s6.7.9的规则。指向通用地址空间的引用和指针的转换可以由任何C++强制转换运算符完成（作为隐式转换）；从通用地址空间到命名地址空间的转换只能使用专用的addrspace_cast运算符来完成。addrspace_cast运算符只能在指针和引用的地址空间之间进行转换，不允许进行其他转换。请注意，__constant和任何其他地址空间之间的转换是不允许的。

// Example assumes generic address space support.
//示例假定支持通用地址空间。
int * genptr; // points to generic address space.
              //指向通用地址空间。

// generic -> named address space conversions.
//泛型->命名地址空间转换。
__private float * ptrfloat = reinterpret_cast<__private float*>(genptr); // illegal.
__private float * ptrfloat = addrspace_cast<__private float*>(genptr); // illegal.
__private int * ptr = addrspace_cast<__private int*>(genptr); // legal.

// named -> generic address space conversion.
//named->通用地址空间转换。
float * genptrfloat = reinterpret_cast<float*>(ptr); // legal.

// disjoint address space conversion.
//不相交的地址空间转换。
__constant int * constptr = addrspace_cast<__constant int*>(genptr); // illegal.

 If generic address space is not supported, any conversion of references/pointers pointing to different address spaces is illegal.

如果不支持通用地址空间，则指向不同地址空间的引用/指针的任何转换都是非法的。

// Example without generic address space support.
//没有通用地址空间支持的示例。
int * privptr; // points to private address space.
               //指向专用地址空间。

// The same address space conversions.
//相同的地址空间转换。
__private float * ptrfloat = reinterpret_cast<__private float*>(privptr); // legal.
__private float * ptrfloat = addrspace_cast<__private float*>(privptr); // illegal.
__private int * ptr = addrspace_cast<__private int*>(privptr); // legal, no op.
float * privptrfloat = reinterpret_cast<float*>(ptr); // legal.

// disjoint address space conversion.
//不相交的地址空间转换。
__constant int * constptr = addrspace_cast<__constant int*>(privptr); // illegal.

3.3.2. References

3.3.2引用

Reference types can be qualified with an address space.

引用类型可以用地址空间限定。

__private int & ref = ...; // references int in __private address space.
                           //引用__private地址空间中的int。

 By default references refer to generic address space objects if generic address space is supported or private address space otherwise, except for dependent types that are not template specializations (see Address space inference).

默认情况下，如果支持通用地址空间，则引用引用通用地址空间对象，否则引用引用私有地址空间，但不是模板专用的依赖类型除外（请参见地址空间推断）。

int & ref = ...; // references int in generic address space if it is
  // supported otherwise in __private address space.
//如果通用地址空间支持引用int，否则引用__private地址空间。

 Address space compatibility checks are performed when references are bound to values. The logic follows the rules from address space pointer conversion (OpenCL C 3.0 s6.7.9).

当引用绑定到值时，将执行地址空间兼容性检查。逻辑遵循地址空间指针转换的规则（OpenCL C 3.0 s6.7.9）。

void f(float &ref, __global float &globref) {
  const int& tmp = ref; // legal - reference to generic/__private address space object
                        // can bind to a temporary object created in __private
                        // address space.
                        //合法的-对通用__private地址空间对象的引用可以绑定到在__private寻址空间中创建的临时对象。

  __global const int& globtmp = globref; // error: reference to global address space
                                         // object cannot bind to a temporary object
                                         // created in __private address space.
                                         //错误：对全局地址空间对象的引用无法绑定到在__private地址空间中创建的临时对象。
}

 3.3.3. Address space inference

3.3.3地址空间推断

This section details what happens if address spaces for types are not provided in the source code explicitly. Most of the logic for address space inference (i.e. default address space) follows rules from OpenCL C 3.0 s6.7.8.

本节详细说明了如果源代码中没有显式提供类型的地址空间会发生什么。地址空间推断的大部分逻辑（即默认地址空间）遵循OpenCL C 3.0 s6.7.8中的规则。

References inherit rules from pointers and therefore refer to generic address space objects by default (see References) if generic address space is supported, otherwise they refer to private address space objects.

引用从指针继承规则，因此如果支持通用地址空间，则默认情况下引用通用地址空间对象（请参阅引用），否则引用专用地址空间对象。

Class static data members are deduced to __global address space. Note, that if the C++ for OpenCL 2021 implementation does not support program scope variables (i.e. __opencl_c_program_scope_global_variables feature from OpenCL C 3.0 s6.1 is unsupported) the address space qualifier must be specified explicitly and it must be the __constant address space.

类静态数据成员被推导为__global地址空间。请注意，如果C++for OpenCL 2021实现不支持程序范围变量（即不支持OpenCL C 3.0 s6.1中的__opencl_c_program_scope_global_variables 功能），则必须显式指定地址空间限定符，并且它必须是__constant地址空间。

All non-static member functions take an implicit object parameter this that is a pointer type. By default the this pointer parameter is in the generic address space if it is supported and in the private address space otherwise. All concrete objects passed as an argument to the implicit this parameter will be converted to this default (generic or private) address space first if such conversion is valid. Therefore, when member functions are called with objects created in disjoint address spaces from the default one, the compilation must fail. To prevent the failure the address space on implicit object parameter this must be specified using address space qualifiers on member functions (see Member function qualifier). For example, use of member functions with objects in __constant address space will always require a __constant member function qualifier as __constant address space is disjoint with any other.

所有非静态成员函数都采用一个隐式对象参数this，该参数是指针类型。默认情况下，如果支持此指针参数，则此指针参数在通用地址空间中，否则此指针参数位于专用地址空间中。如果转换有效，则作为隐式this参数的参数传递的所有具体对象都将首先转换为此默认（通用或专用）地址空间。因此，当使用在与默认地址空间不相交的地址空间中创建的对象调用成员函数时，编译一定会失败。为了防止失败，必须使用成员函数上的地址空间限定符来指定隐式对象参数的地址空间（请参见成员函数限定符）。例如，将成员函数与__constant地址空间中的对象一起使用将始终需要__constant成员函数限定符，因为__constant的地址空间与任何其他地址空间都不相交。

Member function qualifiers can also be used in case address space conversions are undesirable for example for performance reasons. For example, a method can be implemented to exploit memory access coalescing for segments with memory bank.

在不希望进行地址空间转换的情况下，例如出于性能原因，也可以使用成员函数限定符。例如，可以实现一种方法来利用具有内存库的段的内存访问合并。

Address spaces are not deduced for:

以下情况不推导地址空间：

• non-pointer/non-reference template parameters except for template specializations or non-type based template parameters.

•非指针/非引用模板参数，模板专业化或非基于类型的模板参数除外。

• non-pointer/non-reference class members except for static data members that are deduced to the __global address space for C++ for OpenCL 1.0 or C++ for OpenCL 2021 with the __opencl_c_program_scope_global_variables feature.

•非指针/非引用类成员，静态数据成员除外，这些成员被推导到具有__opencl_c_program_scope_global_variables 功能的OpenCL 1.0的C++或OpenCL 2021的C++的__global地址空间。

• non-pointer/non-reference type alias declarations.

•非指针/非引用类型的别名声明。

• decltype expressions.

•decltype表达式。

template <typename T>
void foo() {
  T m; // address space of 'm' will be known at template instantiation time.
       //“m”的地址空间将在模板实例化时已知。
  T * ptr; // 'ptr' points to generic address space object when it is
           // supported otherwise to __private address space.
           //“ptr”指向泛型地址空间对象（如果支持），否则指向__private地址空间。
  T & ref = ...; // 'ref' references an object in generic address space when
                 // it is supported otherwise in __private address space.
                 //“ref”引用通用地址空间中的对象，而__private地址空间中不支持该对象。
};
template <int N>
struct S {
  int i; // 'i' has no address space.
         //“i”没有地址空间。
  static int ii; // 'ii' is in global address space if program scope variables
                 // are supported; otherwise this statement is not legal.
                 //如果支持程序范围变量，则“ii”在全局地址空间中；否则，此声明不合法。
  int * ptr; // 'ptr' points to int in generic address space if it is supported;
             // otherwise to __private address space.
             //“ptr”指向通用地址空间中的int（如果支持）；否则为__private地址空间。
  int & ref = ...; // 'ref' references int in generic address space if it is
                   // supported; otherwise in __private address space.
                   //“ref”在通用地址空间中引用int（如果支持）；否则在__private地址空间中。
};
template <int N>
void bar()
{
  S<N> s; // 's' is in __private address space.
          //“s”在__private地址空间中。
}

 

struct c1 {};
using alias_c1 = c1; // 'alias_c1' is 'c1'.
                     //“alias_c1”是“c1”。
using alias_c1_ptr = c1 *; // 'alias_c1_ptr' is a generic address space pointer to
                           // 'c1' when generic address space is supported; otherwise
                           // it points to 'c1' located in __private address space.
                           //“alias_c1_ptr”是支持通用地址空间时指向“c1”的通用地址空间指针；否则它指向位于__private地址空间中的“c1”。

__kernel void foo()
{
  __local int i;
  decltype(i)* ii; // type of 'ii' is '__local int *__private'.
                   //“ii”的类型为“__local int*__private”。
}

For the placeholder type specifier auto an address space of the outer type is deduced as if it would be any other regular type. However if auto is used in a reference or pointer type, the address space of a pointee is taken from the type of the initialization expression. The logic follows rules for const and volatile qualifiers

对于占位符类型说明符auto，外部类型的地址空间被推断为任何其他常规类型。但是，如果在引用或指针类型中使用auto，则指针对象的地址空间取自初始化表达式的类型。逻辑遵循const和volatile限定符的规则

// This example assumes that generic address space is supported.
//此示例假定支持通用地址空间。
__kernel void foo()
{
  __local int i;
  constexpr int c = 1;

  __constant auto cai = c; // type of 'cai' is '__constant int' (no deduction).
                           //“cai”的类型为“__constant int”（无推论）。

  auto aii = cai; // type of 'aii' is '__private int' (regular deduction).
                  //“aii”的类型为“__private int”（常规扣除）。

  auto *ptr = &i; // type of 'ptr' is '__local int * __private'
                  // (addr space of a pointer is deduced regularly,
                  // addr space of its pointee is taken from 'i').
                  //“ptr”的类型为“__local int*__private”（指针的addr空间是定向推导的，其指针对象的addr空间取自“i”）。

  auto *&refptr = ptr; // type of 'refptr' is '__local int * generic & __private'
                       // (addr space of a reference and type of referencing object
                       // is deduced regularly,
                       // addr space of a pointee is taken from the pointee of 'ptr').
                       //“refptr”的类型为“__local int * generic & __private”（引用的addr空间和引用对象的类型是定向推导的，指针对象的addr空间取自“ptr”中的指针对象）。
}

 3.3.4. Member function qualifier

3.3.4.成员函数限定符

C++ for OpenCL allows specifying an address space qualifier on member functions to signal that they are to be used with objects constructed in a specific address space. This works just the same as qualifying member functions with const or any other qualifiers. The overloading resolution will select the candidate with the most specific address space if multiple candidates are provided. If there is no conversion to an address space among candidates, compilation will fail with a diagnostic.

用于OpenCL的C++允许在成员函数上指定地址空间限定符，以表示它们将与在特定地址空间中构造的对象一起使用。这与使用const或任何其他限定符限定成员函数的作用相同。如果提供了多个候选者，则重载解决方案将选择具有最特定地址空间的候选者。如果候选者之间没有转换到地址空间，则编译将失败并进行诊断。

struct C {
  C() __local {};
  C() __private {};
  constexpr C() __constant {};

  void foo() __local;
  void foo(); // This is implicitly qualified by generic address space
              // if it is supported otherwise by '__private'.
              //如果'__private'以其他方式支持它，则它由通用地址空间隐式限定。
};

__kernel void bar() {
  __local C c1; // will resolve to the first constructor overload.
                //将解析为第一个构造函数重载。
  __private C c2; // will resolve to the second constructor overload.
                  //将解析为第二个构造函数重载。
  __constant C c3{}; // will resolve to the third constructor overload.
                     //将解析为第三个构造函数重载。
  c1.foo(); // will resolve to the first 'foo'.
            //将解析为第一个“foo”。
  c2.foo(); // will resolve to the second 'foo'.
            //将解析为第二个“foo”。
  c3.foo(); // error due to mismatching address spaces - can't convert to
            // '__local' or generic/'__private' address spaces.
            //由于地址空间不匹配而出错-无法转换为'__local'或通用/'__private'地址空间。
}

 All member functions can be qualified by an address space qualifier including constructors and destructors.

所有成员函数都可以由地址空间限定符限定，该限定符包括构造函数和析构函数。

3.3.5. Lambda function

3.3.5.Lambda函数

The address space qualifier can be optionally added for lambda expressions after the attributes. Similar to method qualifiers, they will alter the default address space of lambda call operator that has generic address space by default if it is supported otherwise private address space.

可以为lambda表达式在属性之后添加地址空间限定符。与方法限定符类似，它们将更改默认情况下具有通用地址空间的lambda调用运算符的默认地址空间（如果支持其他专用地址空间）。

__kernel void foo() {
  auto priv1 = []() __private {};
  priv1();
  auto priv2 = []() __global {};
  priv2(); // error: lambda object and its expression have mismatching address space.
           //错误：lambda对象及其表达式的地址空间不匹配。
  __constant auto const3 = []() __constant{};
  const3();

  [&] () __global {} (); // error: lambda temporary is in __private address space.
                         //错误：lambda临时地址在__private地址空间中。

  [&] () mutable __private {} ();
  [&] () __private mutable {} (); // error: mutable specifier should precede address
                                  // space.
                                  //错误：可变说明符应位于地址空间之前。
}

 3.3.6. Implicit special members

3.3.6.隐性特殊成员

The prototype for implicit special members (default, copy or move constructor, copy or move assignment, destructor) has the default address space for an implicit object pointer and reference parameters (see also Member function qualifier). This default address space is generic if it is supported or private otherwise.

隐式特殊成员（默认、复制或移动构造函数、复制或移赋值、析构函数）的原型具有隐式对象指针和引用参数的默认地址空间（另请参见成员函数限定符）。如果支持此默认地址空间，则该地址空间为通用地址空间，否则为专用地址空间。

class C {
  // Has the following implicitly defined member functions.
  //具有以下隐式定义的成员函数。

  // C(); /* implicit 'this' parameter is a pointer to */
          /* object in generic address space if supported, or */
          /* in private address space otherwise. */
          /*隐式“this”参数是指向通用地址空间中对象的指针（如果支持），或者是指向专用地址空间中的对象的指针*/

  // C(const C & par); /* 'this'/'par' is a pointer/reference to */
                       /* object in generic address space */
                       /* if supported, or */
                       /* in private address space otherwise. */
                       /*“this”/“par”是指向通用地址空间中对象的指针/引用（如果支持），或者是指向专用地址空间中的对象的指针或引用（否则）*/

  // C(C && par); /* 'this'/'par' is a pointer/r-val reference to */
                  /* object in generic address space if supported, or */
                  /* in private address space otherwise. */
                  /*“this”/“par”是指向通用地址空间中对象的指针/r-val引用（如果支持），或者在专用地址空间中（否则）*/

  // C & operator=(const C & par); /* 'this'/'par'/return value is */
                                   /* a pointer/reference/reference to */
                                   /* object in generic address space */
                                   /* if supported, or */
                                   /* in private address space otherwise. */
                                   /*“this”/“par”/返回值是指向通用地址空间中对象的指针/引用/引用（如果支持），或者是专用地址空间中的对象*/

  // C & operator=(C && par)'; /* 'this'/'par'/return value is */
                               /* a pointer/r-val reference/reference to */
                               /* object in generic address space, */
                               /* if supported, or */
                               /* in private address space otherwise. */
                               /*“this”/“par”/返回值是指向通用地址空间中对象的指针/r-val引用/引用（如果支持），或者是专用地址空间中的对象的指针/re-val引用*/
};

3.3.7. Builtin operators

3.3.7.内置操作

All builtin operators are available with the specific address spaces, thus no address space conversions (i.e. to generic address space) are performed.

所有内置运算符都可用于特定的地址空间，因此不执行地址空间转换（即到通用地址空间）。

3.3.8. Templates

3.3.8.模板

There is no deduction of address spaces in non-pointer/non-reference template parameters and dependent types (see Address space inference). The address space of a template parameter is deduced during type deduction if it is not explicitly provided in the instantiation.

在非指针/非引用模板参数和依赖类型中没有推导地址空间（请参见地址空间推断）。如果在实例化中没有明确提供模板参数的地址空间，则在类型推导过程中推导该地址空间。

1 template<typename T>
2 void foo(T* i){
3     T var;
4 }
5
6 __global int g;
7 void bar(){
8     foo(&g); // error: template instantiation failed as function scope variable 'var'
9              // appears to be declared in __global address space (see line 3).
               //错误：模板实例化失败，因为函数作用域变量“var”似乎是在__global地址空间中声明的（请参见第3行）。
10 }

It is not legal to specify multiple different address spaces between template definition and instantiation. If multiple different address spaces are specified in a template definition and instantiation, compilation of such a program will fail with a diagnostic. This restriction immediately follows from OpenCL C 3.0 s6.7 that disallows multiple address space qualifiers on a type.

在模板定义和实例化之间指定多个不同的地址空间是不合法的。如果在模板定义和实例化中指定了多个不同的地址空间，则此类程序的编译将失败并进行诊断。这个限制紧跟着OpenCL C 3.0 s6.7，它不允许在一个类型上使用多个地址空间限定符。

template <typename T>
void foo() {
  __private T var;
}

void bar() {
  foo<__global int>(); // error: conflicting address space qualifiers are provided
                       // for 'var', '__global' and '__private'.
                       //错误：为“var”、“__global”和“__private”提供了冲突的地址空间限定符。
}

 Once a template has been instantiated, regular restrictions for address spaces will apply as described in OpenCL C 3.0 s6.7.

一旦模板被实例化，地址空间的常规限制将如OpenCL C 3.0 s6.7中所述。

template<typename T>
void foo(){
  T var;
}

void bar(){
  foo<__global int>(); // error: function scope variable 'var' cannot be declared
                       // in '__global' address space.
                       //错误：函数作用域变量“var”不能在“__global”地址空间中声明。
}

 3.3.9. Temporary materialization

3.3.9.临时具体化

All temporaries are materialized in __private address space. If a reference with another address space is bound to them, a conversion will be generated in case it is valid, otherwise compilation will fail with a diagnostic.

所有临时地址都在__private地址空间中具体化。如果具有另一个地址空间的引用绑定到它们，则将生成转换以防其有效，否则编译将失败并进行诊断。

int bar(const unsigned int &i); // references generic address space object
                                // if generic address space is supported
                                // otherwise private address space object.
                                //如果支持通用地址空间，则引用通用地址空间对象，否则为专用地址空间对象。

void foo() {
  bar(1); // temporary is created in __private address space but (if generic
          // address space is supported) converted to generic address space
          // of parameter reference.
          //临时在__private地址空间中创建，但（如果支持通用地址空间）转换为参数引用的通用地址空间。
}

void f(__global float &ref) {
  __global const int& newref = ref; // error: address space mismatch between
                                    // temporary object created to hold value
                                    // converted float->int and local variable
                                    // (can't convert from __private to __global).
                                    //错误：为保存值转换后的float->int而创建的临时对象和本地变量之间的地址空间不匹配（无法从__private转换为__global）。
}

 3.3.10. Construction, initialization and destruction

3.3.10.构建、初始化和销毁

Construction, initialization and destruction of objects in __private and __global address space follow the general principles of C++. For program scope objects or static objects in the function scope with non-trivial constructors and destructors, the implementation defines an ABI format for runtime initialization and destruction of global objects before/after all kernels are enqueued.

__private和__global地址空间中对象的构造、初始化和销毁遵循C++的一般原则。对于具有非平凡构造函数和析构函数的程序范围对象或函数范围中的静态对象，该实现定义了ABI格式，用于在所有内核入队之前/之后对全局对象进行运行时初始化和销毁。

Objects in __local address space can not have initializers in declarations and therefore a constructor can not be called. All objects created in the local address space have undefined state at the point of their declaration. Developers are free to define a special member function that can initialize local address space objects after their declaration. Any default values provided for the initialization of members in a class declaration are ignored when creating the local address space objects. Since the initialization is performed after the variable declaration, special handling is required for classes with data members that are references because their values can not be overwritten trivially. Destructors of local address space objects are not invoked automatically. They can be called manually.

__local地址空间中的对象在声明中不能有初始值设定项，因此不能调用构造函数。在本地地址空间中创建的所有对象在声明时都有未定义的状态。开发人员可以自由定义一个特殊的成员函数，该函数可以在声明后初始化本地地址空间对象。创建本地地址空间对象时，将忽略为初始化类声明中的成员而提供的任何默认值。由于初始化是在变量声明之后执行的，因此对于具有引用数据成员的类需要特殊处理，因为它们的值不能被平凡地覆盖。本地地址空间对象的析构函数不会自动调用。它们可以手动调用。

class C {
  int m;
// If generic address space is not supported or for performance optimization
// purposes the following members might be required.
//如果不支持通用地址空间或出于性能优化目的，则可能需要以下成员。
public:
  __local C & operator=(const C & par) __local;
  ~C() __local;
};

__kernel void foo() {
  __local C locobj{}; // error: local address space objects can't be initialized
                      //错误：无法初始化本地地址空间对象
  __local C locobj; // uninitialised object.
                    //未初始化的对象。
  locobj = {}; // calling copy assignment operator is allowed.
               //允许调用复制分配运算符。
  locobj.~C(); // local address space object destructors are not invoked
               // automatically.
               //本地地址空间对象析构函数不会自动调用。
}

User defined constructors in __constant address space must be constexpr. Such objects can be initialized using literals and initialization lists if they do not require any user defined conversions.

__constant地址空间中的用户定义构造函数必须是constexpr。如果这些对象不需要任何用户定义的转换，则可以使用文字和初始化列表对其进行初始化。

Objects in __constant address space can be initialized using:

__constant地址空间中的对象可以使用以下方法初始化：

• Literal expressions;

•文字表达式；

• Uniform initialization syntax {};

•统一初始化语法｛｝；

• Using implicit constructors.

•使用隐式构造函数。

• Using constexpr constructors.

•使用constexpr构造函数。

struct C1 {
  int m;
};

struct C2 {
  int m;
  constexpr C2(int init) __constant : m(init) {};
};

__constant C1 c1obj1 = {1};
__constant C1 c1obj2 = C1();
__constant C2 c2obj1(1);

Non-trivial destructors for objects in non-default address spaces (i.e. all other than generic address space when it is supported or __private otherwise) are not required to be supported by implementations. The macro __opencl_cpp_destructor_with_address_spaces, which is defined if and only if such destructors are supported by an implementation, can be used to check whether this functionality can be used in kernel sources. Additionally destructors with global objects might not be supported even if address spaces are supported with destructors in general. Such functionality is indicated by the presence of the __opencl_cpp_global_destructor macro. If the macro __opencl_cpp_global_destructor is defined then __opencl_cpp_destructor_with_address_spaces must also be defined.

非默认地址空间中对象的非平凡析构函数（即，支持通用地址空间或__private地址空间以外的所有对象）不需要由实现支持。宏__opencl_cpp_destructor_with_address_spaces可用于检查是否可以在内核源中使用此功能，该宏是在实现支持此类析构函数时定义的。此外，即使析构函数通常支持地址空间，也可能不支持具有全局对象的析构函数。这种功能通过__opencl_cpp_global_destructor宏的存在来指示。如果定义了宏__opencl_cpp_global_destructor，则还必须定义__opencl.cpp_destructor_with_address_spaces。

Note that the destruction of objects in named address spaces __global, __local, or __private can be performed using destructors with default address space (i.e. generic) by utilising address space conversions

请注意，命名地址空间__global、__local或__private中的对象的销毁可以通过使用地址空间转换，使用具有默认地址空间（即通用型）的析构函数来执行

1 // Example assumes generic address space support.
  //示例假定支持通用地址空间。
2 class C {
3 public:
4 #ifdef __opencl_cpp_destructor_with_address_spaces
5     ~C() __local;
6 #else
7     ~C();
8 #endif
9 };
10
11 kernel void foo() {
12 __local C locobj;
13 locobj.~C(); // uses destructor in local address space (on line 5)
14              // if such destructors are supported,
15              // otherwise uses generic address space destructor (on line 7)
16              // converting to generic address prior to call into destructor.
                //如果支持此类析构函数，则在本地地址空间（第5行）中使用析构函数；否则，在调用析构函数之前，使用通用地址空间析构函数（第7行）转换为通用地址。
17 }

However, when generic address space feature is unsupported, absence of destructor support with address spaces results in a compilation failure when such destructors overloaded with non-default address spaces are encountered in the kernel code.

然而，当通用地址空间功能不受支持时，如果在内核代码中遇到用非默认地址空间重载的析构函数，则缺少对地址空间的析构构函数支持会导致编译失败。

// Example assumes generic address space is not supported.
//示例假定不支持通用地址空间。
class C {
public:
  ~C();
};

kernel void foo() {
  __local C locobj;
  locobj.~C(); // error due to illegal conversion of 'this' from __local
               // to __private address space pointer.
               //由于“this”从__local到__private地址空间指针的非法转换而导致错误。
}

3.3.11. Nested pointers

3.3.11.嵌套指针

C++ for OpenCL does not allow implicit address space conversions in nested pointers even with compatible address spaces. The following rules apply when converting between address spaces in nested pointers:

OpenCL的C++不允许在嵌套指针中进行隐式地址空间转换，即使使用兼容的地址空间也是如此。在嵌套指针中的地址空间之间进行转换时，适用以下规则：

• Implicit conversions of address spaces in nested pointers are disallowed.

•不允许对嵌套指针中的地址空间进行隐式转换。

• Any address space conversion in nested pointers with safe casts (e.g. const_cast, static_cast, addrspace_cast) is disallowed.

•不允许在具有安全强制转换的嵌套指针（例如const_cast、static_cast、addrspace_cast）中进行任何地址空间转换。

• Any address space conversion in nested pointers can be done using low level C-style or reinterpret_cast. No compatibility check is performed for address spaces in nested pointers.

•嵌套指针中的任何地址空间转换都可以使用低级C样式或interpret_cast来完成。不对嵌套指针中的地址空间执行兼容性检查。

local int * * locdefptr;
constant int * * cnstdefptr;
int * * defdefptr;
defdefptr = const_cast<int * *>(locdefptr); // illegal.
defdefptr = static_cast<int * *>(cnstdefptr); // illegal.
defdefptr = addrspace_cast<int * *>(cnstdefptr); // illegal.
defdefptr = reinterpret_cast<int * *>(locdefptr); // legal.
defdefptr = reinterpret_cast<int * *>(cnstdefptr); // legal.

 3.3.12. Address space removal type trait

3.3.12.地址空间移除类型特征

template struct __remove_address_space;

C++ for OpenCL 2021 supports the type trait __remove_address_space that provides the member typedef type which is the same as T, except that its topmost address space qualifier is removed. Its effect is analogous to remove_const and other similar type traits in C++17 [meta.trans]. The trait only removes an address space qualifier from a given type, therefore, all other type qualifiers such as const or volatile remain unchanged.

OpenCL 2021的C++支持类型特征 __remove_addre_space，该类型特征提供与T相同的成员typedef类型，只是删除了其最顶层的地址空间限定符。它的效果类似于C++17[meta.trans]中的remove_const和其他类似的类型特征。该特征仅从给定类型中删除地址空间限定符，因此，所有其他类型限定符（如const或volatile）保持不变。

template<typename T>
void foo(T *par) {
  T var1; // error: function scope variable cannot be declared in global
          // address space.
          //错误：函数作用域变量不能在全局地址空间中声明。
  __private T var2; // error: conflicting address space qualifiers are provided
                    // between types '__private T' and '__global int'.
                    //错误：在类型“__private T”和“__global int”之间提供了冲突的地址空间限定符。
  __private typename __remove_address_space<T>::type var3; // type of var3 is __private int.
                                                           //var3的类型是__private int。
}

void bar() {
  __global int* ptr;
  foo(ptr);
}

 3.4. C++ casts

3.4.C++类型强制转换

C++ has three cast operators in addition to C-style casts. Additional logic specific to address spaces are applied to all casts as detailed in conversions with address spaces. reinterpret_cast has some additional functionality:

除了C样式的强制转换外，C++还有三个强制转换运算符。特定于地址空间的附加逻辑将应用于所有强制转换，如使用地址空间的转换中所述。interpret_cast具有一些附加功能：

• Conversion between vectors and scalars are allowed.

•允许矢量和标量之间的转换。

• Conversion between OpenCL types are disallowed.

•不允许在OpenCL类型之间进行转换。

3.4.1. Vectors and scalars

3.4.1.矢量和标量

reinterpret_cast reinterprets between integral types like integers and pointers. In C++ for openCL this also includes vector types, and so using reinterpret_cast between vectors and scalars is also possible, as long as the size of the vectors are the same.

interpret_cast在整数和指针等积分类型之间进行重新解释。在openCL的C++中，这也包括向量类型，因此只要向量的大小相同，在向量和标量之间使用reinterpret_cast也是可能的。

int i;
short2 s2 = reinterpret_cast<short2>(i); // legal.
int2 i2 = reinterpret_cast<int2>(i); // illegal.

short8 s8;
int4 i4 = reinterpret_cast<int4>(s8); // legal.
long l4 = reinterpret_cast<long>(s8); // illegal.

 3.4.2. OpenCL types

3.4.2.OpenCL类型

Some of the OpenCL-specific types, defined as "Other Built-in Data Types" in OpenCL C 3.0 s6.3.3, are convertible to integer literals, but since they are not conceptually integral, they can not be used with reinterpret_cast. Therefore conversions of an OpenCL-specific type to any distinct type are illegal.

一些OpenCL特定的类型，在OpenCL C 3.0 s6.3.3中被定义为“其他内置数据类型”，可以转换为整数文字，但由于它们在概念上不是整数，因此不能与interpret_cast一起使用。因此，将OpenCL特定类型转换为任何不同类型都是非法的。

queue_t q;
reserve_id_t id = reinterpret_cast<reserve_id_t>(q); // illegal.
int i = reinterpret_cast<int>(id); // illegal.

 3.5. Kernel functions

3.5.内核函数

Kernel functions have implicit C linkage (C++17 [dcl.link]) which means that C++ specific features are not supported. Therefore, the kernel functions:

内核函数具有隐含的C链接（C++17[dcl.link]），这意味着不支持C++特定的功能。因此，内核的功能是：

• Can not be class members (C++17 [class.mfct]);

•不能是类成员（C++17[class.mfct]）；

• Can not be overloaded (C++17 [over]);

•不能过载（C++17[over]）；

• Can not be function templates (C++17 [temp.fct]).

•不能是函数模板（C++17[temp.fct]）。

Moreover the types used in parameters of the kernel functions must be:

此外，内核函数的参数中使用的类型必须是：

• Trivial and standard-layout types C++17 [basic.types] (plain old data types) for parameters passed by value;

•按值传递的参数的琐碎和标准布局类型C++17[基本类型]（普通的旧数据类型）；

• Standard-layout types for pointer parameters. The same applies to references[2] if an implementation supports them in kernel parameters.

•指针参数的标准布局类型。如果实现在内核参数中支持引用[2]，那么引用[2]也是如此。

These are additional restrictions to the list detailed in OpenCL C 3.0 s6.11.

这些是对OpenCL C 3.0 s6.11中详细列出的列表的附加限制。

[1] The macro belongs to the list of C++20’s feature test macros

[1] 该宏属于C++20的功能测试宏列表

[2] Whether C++ features (e.g references) can be used in functions with C linkage is implementation-defined (C++17 [dcl.link]).

[2] C++特性（例如引用）是否可以在具有C链接的函数中使用是由实现定义的（C++17[dcl.link]）。

Chapter 4. Normative References

第4章 规范性引用文件

1. “The OpenCL Specification, Version 3.0”, https://www.khronos.org/registry/OpenCL/.

1.“OpenCL规范，3.0版”，https://www.khronos.org/registry/OpenCL/.

2. “The OpenCL C Specification, Version 3.0”, https://www.khronos.org/registry/OpenCL/.

2.“OpenCL C规范3.0版”，https://www.khronos.org/registry/OpenCL/.

3. “ISO/IEC 14882:2017 - Programming languages — C++”, https://www.iso.org/standard/ 68564.html. References are to sections of this specific version, referred to as the “The C++17 Specification”, although other versions exist.

3.“ISO/IEC 14882:2017-程序设计语言C++”，https://www.iso.org/standard/68564.html.参考是指该特定版本的部分，称为“C++17规范”，尽管存在其他版本。

4. “ISO/IEC TR 18037:2008 Programming languages - C - Extensions to support embedded processors”, https://www.iso.org/standard/51126.html. References are to sections of this specific version, referred to as the “The Embedded C Specification”, although other versions exist.

4.“ISO/IEC TR 18037:2008程序设计语言-C-支持嵌入式处理器的扩展”，https://www.iso.org/standard/51126.html.参考是指该特定版本的章节，称为“嵌入式C规范”，尽管存在其他版本。

Acknowledgements

致谢

The C++ for OpenCL documentation is the result of the contributions of many people. Following is a partial list of the contributors, including the company that they represented at the time of their contribution:

C++for OpenCL文档是许多人贡献的结果。以下是出资人的部分名单，包括他们出资时所代表的公司：

• Anastasia Stulova, Arm

• Neil Hickey, Arm

• Sven van Haastregt, Arm

• Marco Antognini, Arm

• Kevin Petit, Arm

• Stuart Brady, Arm

• Ole Strøhm, Arm

• Justas Janickas, Arm