r318601 - [OpenMP] Show error if VLAs are not supported
Jonas Hahnfeld via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 20 16:39:53 PST 2017
I'm sorry, I didn't get a notification for that failure. I'll reply to
the Phabricator.
Am 2017-11-20 19:37, schrieb Artem Belevich:
> Proposed fix: https://reviews.llvm.org/D40275
>
> On Mon, Nov 20, 2017 at 4:13 PM, Artem Belevich <tra at google.com>
> wrote:
>
>> This change breaks CUDA as clang now reports an error during
>> device-side compilation when VLA is used in the *host-side* code.
>>
> http://lab.llvm.org:8011/builders/clang-cuda-build/builds/15591/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio
>> [13]
>>
>> E.g. I would expect this code to compile successfully, producing
>> empty device-side binary:
>>
>> void host_func(int i) {
>> int vla[i];
>> }
>>
>> However it currently fails:
>> #bin/clang++ --cuda-device-only --cuda-gpu-arch=sm_35 -o vla.o
>> vla.cu [14]
>>
>> vla.cu:4:10: error: variable length arrays are not supported for the
>> current target
>> int vla[i];
>> ^
>> 1 error generated when compiling for sm_35.
>>
>> On Sat, Nov 18, 2017 at 1:00 PM, Jonas Hahnfeld via cfe-commits
>> <cfe-commits at lists.llvm.org> wrote:
>>
>>> Author: hahnfeld
>>> Date: Sat Nov 18 13:00:46 2017
>>> New Revision: 318601
>>>
>>> URL: http://llvm.org/viewvc/llvm-project?rev=318601&view=rev [1]
>>> Log:
>>> [OpenMP] Show error if VLAs are not supported
>>>
>>> Some target devices (e.g. Nvidia GPUs) don't support dynamic stack
>>> allocation and hence no VLAs. Print errors with description
>>> instead
>>> of failing in the backend or generating code that doesn't work.
>>>
>>> This patch handles explicit uses of VLAs (local variable in target
>>> or declare target region) or implicitly generated (private) VLAs
>>> for reductions on VLAs or on array sections with non-constant
>>> size.
>>>
>>> Differential Revision: https://reviews.llvm.org/D39505 [2]
>>>
>>> Added:
>>> cfe/trunk/test/OpenMP/target_vla_messages.cpp
>>> Modified:
>>> cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
>>> cfe/trunk/include/clang/Basic/TargetInfo.h
>>> cfe/trunk/include/clang/Sema/Sema.h
>>> cfe/trunk/lib/Basic/TargetInfo.cpp
>>> cfe/trunk/lib/Basic/Targets/NVPTX.cpp
>>> cfe/trunk/lib/Basic/Targets/SPIR.h
>>> cfe/trunk/lib/Sema/SemaOpenMP.cpp
>>> cfe/trunk/lib/Sema/SemaType.cpp
>>>
>>> Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=318601&r1=318600&r2=318601&view=diff
>>> [3]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
>>> (original)
>>> +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Sat Nov
>>> 18 13:00:46 2017
>>> @@ -141,6 +141,10 @@ def err_vla_decl_has_extern_linkage : Er
>>> "variable length array declaration cannot have 'extern'
>>> linkage">;
>>> def ext_vla_folded_to_constant : Extension<
>>> "variable length array folded to constant array as an
>>> extension">, InGroup<GNUFoldingConstant>;
>>> +def err_vla_unsupported : Error<
>>> + "variable length arrays are not supported for the current
>>> target">;
>>> +def note_vla_unsupported : Note<
>>> + "variable length arrays are not supported for the current
>>> target">;
>>>
>>> // C99 variably modified types
>>> def err_variably_modified_template_arg : Error<
>>> @@ -8985,6 +8989,8 @@ def err_omp_reduction_non_addressable_ex
>>> "expected addressable reduction item for the task-based
>>> directives">;
>>> def err_omp_reduction_with_nogroup : Error<
>>> "'reduction' clause cannot be used with 'nogroup' clause">;
>>> +def err_omp_reduction_vla_unsupported : Error<
>>> + "cannot generate code for reduction on %select{|array section,
>>> which requires a }0variable length array">;
>>> } // end of OpenMP category
>>>
>>> let CategoryName = "Related Result Type Issue" in {
>>>
>>> Modified: cfe/trunk/include/clang/Basic/TargetInfo.h
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=318601&r1=318600&r2=318601&view=diff
>>> [4]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/include/clang/Basic/TargetInfo.h (original)
>>> +++ cfe/trunk/include/clang/Basic/TargetInfo.h Sat Nov 18 13:00:46
>>> 2017
>>> @@ -60,6 +60,7 @@ protected:
>>> // values are specified by the TargetInfo constructor.
>>> bool BigEndian;
>>> bool TLSSupported;
>>> + bool VLASupported;
>>> bool NoAsmVariants; // True if {|} are normal characters.
>>> bool HasFloat128;
>>> unsigned char PointerWidth, PointerAlign;
>>> @@ -939,6 +940,9 @@ public:
>>> return MaxTLSAlign;
>>> }
>>>
>>> + /// \brief Whether target supports variable-length arrays.
>>> + bool isVLASupported() const { return VLASupported; }
>>> +
>>> /// \brief Whether the target supports SEH __try.
>>> bool isSEHTrySupported() const {
>>> return getTriple().isOSWindows() &&
>>>
>>> Modified: cfe/trunk/include/clang/Sema/Sema.h
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=318601&r1=318600&r2=318601&view=diff
>>> [5]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/include/clang/Sema/Sema.h (original)
>>> +++ cfe/trunk/include/clang/Sema/Sema.h Sat Nov 18 13:00:46 2017
>>> @@ -8653,10 +8653,18 @@ public:
>>> NamedDeclSetType
>>> &SameDirectiveDecls);
>>> /// Check declaration inside target region.
>>> void checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D);
>>> - /// Return true inside OpenMP target region.
>>> + /// Return true inside OpenMP declare target region.
>>> bool isInOpenMPDeclareTargetContext() const {
>>> return IsInOpenMPDeclareTargetContext;
>>> }
>>> + /// Return true inside OpenMP target region.
>>> + bool isInOpenMPTargetExecutionDirective() const;
>>> + /// Return true if (un)supported features for the current
>>> target should be
>>> + /// diagnosed if OpenMP (offloading) is enabled.
>>> + bool shouldDiagnoseTargetSupportFromOpenMP() const {
>>> + return !getLangOpts().OpenMPIsDevice ||
>>> isInOpenMPDeclareTargetContext() ||
>>> + isInOpenMPTargetExecutionDirective();
>>> + }
>>>
>>> /// Return the number of captured regions created for an OpenMP
>>> directive.
>>> static int getOpenMPCaptureLevels(OpenMPDirectiveKind Kind);
>>>
>>> Modified: cfe/trunk/lib/Basic/TargetInfo.cpp
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=318601&r1=318600&r2=318601&view=diff
>>> [6]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/lib/Basic/TargetInfo.cpp (original)
>>> +++ cfe/trunk/lib/Basic/TargetInfo.cpp Sat Nov 18 13:00:46 2017
>>> @@ -31,6 +31,7 @@ TargetInfo::TargetInfo(const llvm::Tripl
>>> // SPARC. These should be overridden by concrete targets as
>>> needed.
>>> BigEndian = !T.isLittleEndian();
>>> TLSSupported = true;
>>> + VLASupported = true;
>>> NoAsmVariants = false;
>>> HasFloat128 = false;
>>> PointerWidth = PointerAlign = 32;
>>>
>>> Modified: cfe/trunk/lib/Basic/Targets/NVPTX.cpp
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.cpp?rev=318601&r1=318600&r2=318601&view=diff
>>> [7]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/lib/Basic/Targets/NVPTX.cpp (original)
>>> +++ cfe/trunk/lib/Basic/Targets/NVPTX.cpp Sat Nov 18 13:00:46 2017
>>> @@ -41,6 +41,7 @@ NVPTXTargetInfo::NVPTXTargetInfo(const l
>>> "NVPTX only supports 32- and 64-bit modes.");
>>>
>>> TLSSupported = false;
>>> + VLASupported = false;
>>> AddrSpaceMap = &NVPTXAddrSpaceMap;
>>> UseAddrSpaceMapMangling = true;
>>>
>>> Modified: cfe/trunk/lib/Basic/Targets/SPIR.h
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/SPIR.h?rev=318601&r1=318600&r2=318601&view=diff
>>> [8]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/lib/Basic/Targets/SPIR.h (original)
>>> +++ cfe/trunk/lib/Basic/Targets/SPIR.h Sat Nov 18 13:00:46 2017
>>> @@ -43,6 +43,7 @@ public:
>>> assert(getTriple().getEnvironment() ==
>>> llvm::Triple::UnknownEnvironment &&
>>> "SPIR target must use unknown environment type");
>>> TLSSupported = false;
>>> + VLASupported = false;
>>> LongWidth = LongAlign = 64;
>>> AddrSpaceMap = &SPIRAddrSpaceMap;
>>> UseAddrSpaceMapMangling = true;
>>>
>>> Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=318601&r1=318600&r2=318601&view=diff
>>> [9]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
>>> +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Sat Nov 18 13:00:46 2017
>>> @@ -1303,6 +1303,17 @@ unsigned Sema::getOpenMPNestingLevel() c
>>> return DSAStack->getNestingLevel();
>>> }
>>>
>>> +bool Sema::isInOpenMPTargetExecutionDirective() const {
>>> + return
>>> (isOpenMPTargetExecutionDirective(DSAStack->getCurrentDirective())
>>> &&
>>> + !DSAStack->isClauseParsingMode()) ||
>>> + DSAStack->hasDirective(
>>> + [](OpenMPDirectiveKind K, const DeclarationNameInfo
>>> &,
>>> + SourceLocation) -> bool {
>>> + return isOpenMPTargetExecutionDirective(K);
>>> + },
>>> + false);
>>> +}
>>> +
>>> VarDecl *Sema::IsOpenMPCapturedDecl(ValueDecl *D) {
>>> assert(LangOpts.OpenMP && "OpenMP is not allowed");
>>> D = getCanonicalDecl(D);
>>> @@ -1315,18 +1326,8 @@ VarDecl *Sema::IsOpenMPCapturedDecl(Valu
>>> // inserted here once support for 'declare target' is added.
>>> //
>>> auto *VD = dyn_cast<VarDecl>(D);
>>> - if (VD && !VD->hasLocalStorage()) {
>>> - if
>>> (isOpenMPTargetExecutionDirective(DSAStack->getCurrentDirective())
>>> &&
>>> - !DSAStack->isClauseParsingMode())
>>> - return VD;
>>> - if (DSAStack->hasDirective(
>>> - [](OpenMPDirectiveKind K, const DeclarationNameInfo
>>> &,
>>> - SourceLocation) -> bool {
>>> - return isOpenMPTargetExecutionDirective(K);
>>> - },
>>> - false))
>>> - return VD;
>>> - }
>>> + if (VD && !VD->hasLocalStorage() &&
>>> isInOpenMPTargetExecutionDirective())
>>> + return VD;
>>>
>>> if (DSAStack->getCurrentDirective() != OMPD_unknown &&
>>> (!DSAStack->isClauseParsingMode() ||
>>> @@ -9812,6 +9813,12 @@ static bool ActOnOMPReductionKindClause(
>>> if ((OASE && !ConstantLengthOASE) ||
>>> (!OASE && !ASE &&
>>>
>>> D->getType().getNonReferenceType()->isVariablyModifiedType())) {
>>> + if (!Context.getTargetInfo().isVLASupported() &&
>>> + S.shouldDiagnoseTargetSupportFromOpenMP()) {
>>> + S.Diag(ELoc, diag::err_omp_reduction_vla_unsupported) <<
>>> !!OASE;
>>> + S.Diag(ELoc, diag::note_vla_unsupported);
>>> + continue;
>>> + }
>>> // For arrays/array sections only:
>>> // Create pseudo array type for private copy. The size for
>>> this array will
>>> // be generated during codegen.
>>>
>>> Modified: cfe/trunk/lib/Sema/SemaType.cpp
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=318601&r1=318600&r2=318601&view=diff
>>> [10]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/lib/Sema/SemaType.cpp (original)
>>> +++ cfe/trunk/lib/Sema/SemaType.cpp Sat Nov 18 13:00:46 2017
>>> @@ -2183,6 +2183,12 @@ QualType Sema::BuildArrayType(QualType T
>>> // CUDA device code doesn't support VLAs.
>>> if (getLangOpts().CUDA && T->isVariableArrayType())
>>> CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) <<
>>> CurrentCUDATarget();
>>> + // Some targets don't support VLAs.
>>> + if (T->isVariableArrayType() &&
>>> !Context.getTargetInfo().isVLASupported() &&
>>> + shouldDiagnoseTargetSupportFromOpenMP()) {
>>> + Diag(Loc, diag::err_vla_unsupported);
>>> + return QualType();
>>> + }
>>>
>>> // If this is not C99, extwarn about VLA's and C99 array size
>>> modifiers.
>>> if (!getLangOpts().C99) {
>>>
>>> Added: cfe/trunk/test/OpenMP/target_vla_messages.cpp
>>> URL:
>>>
>>
> http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=318601&view=auto
>>> [11]
>>>
>>
> ==============================================================================
>>> --- cfe/trunk/test/OpenMP/target_vla_messages.cpp (added)
>>> +++ cfe/trunk/test/OpenMP/target_vla_messages.cpp Sat Nov 18
>>> 13:00:46 2017
>>> @@ -0,0 +1,201 @@
>>> +// PowerPC supports VLAs.
>>> +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple
>>> powerpc64le-unknown-unknown
>>> -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o
>>> %t-ppc-host-ppc.bc
>>> +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple
>>> powerpc64le-unknown-unknown
>>> -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm %s
>>> -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-ppc.bc
>>> -o %t-ppc-device.ll
>>> +
>>> +// Nvidia GPUs don't support VLAs.
>>> +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple
>>> powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda
>>> -emit-llvm-bc %s -o %t-ppc-host-nvptx.bc
>>> +// RUN: %clang_cc1 -verify -DNO_VLA -fopenmp -x c++ -triple
>>> nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda
>>> -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path
>>> %t-ppc-host-nvptx.bc -o %t-nvptx-device.ll
>>> +
>>> +#ifndef NO_VLA
>>> +// expected-no-diagnostics
>>> +#endif
>>> +
>>> +#pragma omp declare target
>>> +void declare(int arg) {
>>> + int a[2];
>>> +#ifdef NO_VLA
>>> + // expected-error at +2 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> + int vla[arg];
>>> +}
>>> +
>>> +void declare_parallel_reduction(int arg) {
>>> + int a[2];
>>> +
>>> +#pragma omp parallel reduction(+: a)
>>> + { }
>>> +
>>> +#pragma omp parallel reduction(+: a[0:2])
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +3 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +2 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp parallel reduction(+: a[0:arg])
>>> + { }
>>> +}
>>> +#pragma omp end declare target
>>> +
>>> +template <typename T>
>>> +void target_template(int arg) {
>>> +#pragma omp target
>>> + {
>>> +#ifdef NO_VLA
>>> + // expected-error at +2 {{variable length arrays are not
>>> supported for the current target}}
>>> +#endif
>>> + T vla[arg];
>>> + }
>>> +}
>>> +
>>> +void target(int arg) {
>>> +#pragma omp target
>>> + {
>>> +#ifdef NO_VLA
>>> + // expected-error at +2 {{variable length arrays are not
>>> supported for the current target}}
>>> +#endif
>>> + int vla[arg];
>>> + }
>>> +
>>> +#pragma omp target
>>> + {
>>> +#pragma omp parallel
>>> + {
>>> +#ifdef NO_VLA
>>> + // expected-error at +2 {{variable length arrays are not
>>> supported for the current target}}
>>> +#endif
>>> + int vla[arg];
>>> + }
>>> + }
>>> +
>>> + target_template<long>(arg);
>>> +}
>>> +
>>> +void teams_reduction(int arg) {
>>> + int a[2];
>>> + int vla[arg];
>>> +
>>> +#pragma omp target map(a)
>>> +#pragma omp teams reduction(+: a)
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +4 {{cannot generate code for reduction on
>>> variable length array}}
>>> + // expected-note at +3 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(vla)
>>> +#pragma omp teams reduction(+: vla)
>>> + { }
>>> +
>>> +#pragma omp target map(a[0:2])
>>> +#pragma omp teams reduction(+: a[0:2])
>>> + { }
>>> +
>>> +#pragma omp target map(vla[0:2])
>>> +#pragma omp teams reduction(+: vla[0:2])
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +4 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +3 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(a[0:arg])
>>> +#pragma omp teams reduction(+: a[0:arg])
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +4 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +3 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(vla[0:arg])
>>> +#pragma omp teams reduction(+: vla[0:arg])
>>> + { }
>>> +}
>>> +
>>> +void parallel_reduction(int arg) {
>>> + int a[2];
>>> + int vla[arg];
>>> +
>>> +#pragma omp target map(a)
>>> +#pragma omp parallel reduction(+: a)
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +4 {{cannot generate code for reduction on
>>> variable length array}}
>>> + // expected-note at +3 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(vla)
>>> +#pragma omp parallel reduction(+: vla)
>>> + { }
>>> +
>>> +#pragma omp target map(a[0:2])
>>> +#pragma omp parallel reduction(+: a[0:2])
>>> + { }
>>> +
>>> +#pragma omp target map(vla[0:2])
>>> +#pragma omp parallel reduction(+: vla[0:2])
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +4 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +3 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(a[0:arg])
>>> +#pragma omp parallel reduction(+: a[0:arg])
>>> + { }
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +4 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +3 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(vla[0:arg])
>>> +#pragma omp parallel reduction(+: vla[0:arg])
>>> + { }
>>> +}
>>> +
>>> +void for_reduction(int arg) {
>>> + int a[2];
>>> + int vla[arg];
>>> +
>>> +#pragma omp target map(a)
>>> +#pragma omp parallel
>>> +#pragma omp for reduction(+: a)
>>> + for (int i = 0; i < arg; i++) ;
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +5 {{cannot generate code for reduction on
>>> variable length array}}
>>> + // expected-note at +4 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(vla)
>>> +#pragma omp parallel
>>> +#pragma omp for reduction(+: vla)
>>> + for (int i = 0; i < arg; i++) ;
>>> +
>>> +#pragma omp target map(a[0:2])
>>> +#pragma omp parallel
>>> +#pragma omp for reduction(+: a[0:2])
>>> + for (int i = 0; i < arg; i++) ;
>>> +
>>> +#pragma omp target map(vla[0:2])
>>> +#pragma omp parallel
>>> +#pragma omp for reduction(+: vla[0:2])
>>> + for (int i = 0; i < arg; i++) ;
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +5 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +4 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(a[0:arg])
>>> +#pragma omp parallel
>>> +#pragma omp for reduction(+: a[0:arg])
>>> + for (int i = 0; i < arg; i++) ;
>>> +
>>> +#ifdef NO_VLA
>>> + // expected-error at +5 {{cannot generate code for reduction on
>>> array section, which requires a variable length array}}
>>> + // expected-note at +4 {{variable length arrays are not supported
>>> for the current target}}
>>> +#endif
>>> +#pragma omp target map(vla[0:arg])
>>> +#pragma omp parallel
>>> +#pragma omp for reduction(+: vla[0:arg])
>>> + for (int i = 0; i < arg; i++) ;
>>> +}
>>>
>>> _______________________________________________
>>> cfe-commits mailing list
>>> cfe-commits at lists.llvm.org
>>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits [12]
>>
>> --
>>
>> --Artem Belevich
>
> --
>
> --Artem Belevich
>
> Links:
> ------
> [1] http://llvm.org/viewvc/llvm-project?rev=318601&view=rev
> [2] https://reviews.llvm.org/D39505
> [3]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=318601&r1=318600&r2=318601&view=diff
> [4]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=318601&r1=318600&r2=318601&view=diff
> [5]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=318601&r1=318600&r2=318601&view=diff
> [6]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=318601&r1=318600&r2=318601&view=diff
> [7]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.cpp?rev=318601&r1=318600&r2=318601&view=diff
> [8]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/SPIR.h?rev=318601&r1=318600&r2=318601&view=diff
> [9]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=318601&r1=318600&r2=318601&view=diff
> [10]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=318601&r1=318600&r2=318601&view=diff
> [11]
> http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=318601&view=auto
> [12] http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
> [13]
> http://lab.llvm.org:8011/builders/clang-cuda-build/builds/15591/steps/ninja%20build%20simple%20CUDA%20tests/logs/stdio
> [14] http://vla.cu
More information about the cfe-commits
mailing list