Skip to content

Commit 6f0ad1a

Browse files
mdtoguchimdtoguchi
authored
[Driver][SYCL] Provide ability to use an external host compiler (#3530)
The default toolchain when performing -fsycl compilations is to use the clang compiler for the device and host compilations. Enable the ability to use an external host compiler (e.g. gcc/g++) to be used when creating the host objects. This introduces 2 options: -fsycl-host-compiler=<arg> When specified, this option informs the compiler driver that the host compilation step that is performed as part of the greater compilation flow will be performed by the compiler <arg>. It is expected that <arg> is the compiler to be called, either by name (in which the PATH will be used to discover it) or a fully qualified directory with compiler to invoke. This option is only useful when used with when '-fsycl' is provided on the command line. -fsycl-host-compiler-options="opts" Passes along the space separated quoted "opts" string as option arguments to the compiler specified with the -fsycl-host-compiler=<arg> option. It is expected that the options used here are compatible with the compiler specified via -fsycl-host-compiler=<arg>. This option does not have any affect without -fsycl-host-compiler=<arg>.
1 parent ff942e9 commit 6f0ad1a

File tree

9 files changed

+363
-2
lines changed

9 files changed

+363
-2
lines changed

‎clang/include/clang/Basic/DiagnosticDriverKinds.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -106,6 +106,8 @@ def err_drv_use_of_Z_option : Error<
106106
"unsupported use of internal gcc -Z option '%0'">;
107107
def err_drv_output_argument_with_multiple_files : Error<
108108
"cannot specify -o when generating multiple output files">;
109+
def err_drv_output_type_with_host_compiler : Error<
110+
"unsupported output type when using external host compiler">;
109111
def err_drv_out_file_argument_with_multiple_sources : Error<
110112
"cannot specify '%0%1' when compiling multiple source files">;
111113
def err_no_external_assembler : Error<

‎clang/include/clang/Driver/Options.td

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2524,6 +2524,13 @@ def fsycl_help : Flag<["-"], "fsycl-help">, Alias<fsycl_help_EQ>,
25242524
def fsycl_libspirv_path_EQ : Joined<["-"], "fsycl-libspirv-path=">,
25252525
Flags<[CC1Option, CoreOption]>, HelpText<"Path to libspirv library">;
25262526
def fno_sycl_libspirv : Flag<["-"], "fno-sycl-libspirv">, HelpText<"Disable check for libspirv">;
2527+
def fsycl_host_compiler_EQ : Joined<["-"], "fsycl-host-compiler=">,
2528+
Flags<[CoreOption]>, HelpText<"Specify C++ compiler binary to perform host "
2529+
"compilation with during SYCL offload compiles.">;
2530+
def fsycl_host_compiler_options_EQ : Joined<["-"], "fsycl-host-compiler-options=">,
2531+
Flags<[CoreOption]>, HelpText<"When performing the host compilation with "
2532+
"-fsycl-host-compiler specified, use the given options during that compile. "
2533+
"Options are expected to be a quoted list of space separated options.">;
25272534
def fsyntax_only : Flag<["-"], "fsyntax-only">,
25282535
Flags<[NoXarchOption,CoreOption,CC1Option,FC1Option]>, Group<Action_Group>;
25292536
def ftabstop_EQ : Joined<["-"], "ftabstop=">, Group<f_Group>;

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 170 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,7 @@
4040
#include "llvm/Option/ArgList.h"
4141
#include "llvm/Support/Casting.h"
4242
#include "llvm/Support/CodeGen.h"
43+
#include "llvm/Support/CommandLine.h"
4344
#include "llvm/Support/Compiler.h"
4445
#include "llvm/Support/Compression.h"
4546
#include "llvm/Support/FileSystem.h"
@@ -4195,6 +4196,167 @@ static bool ContainsWrapperAction(const Action *A) {
41954196
return false;
41964197
}
41974198

4199+
// Put together an external compiler compilation call which is used instead
4200+
// of the clang invocation for the host compile of an offload compilation.
4201+
// Enabling command line: clang++ -fsycl -fsycl-host-compiler=<HostExe>
4202+
// <ClangOpts> -fsycl-host-compiler-options=<HostOpts>
4203+
// Any <ClangOpts> used which are phase limiting (preprocessing, assembly,
4204+
// object generation) are specifically handled here by specifying the
4205+
// equivalent phase limiting option(s).
4206+
// It is expected that any user <HostOpts> options passed will be placed
4207+
// after any implied options set here. This will have overriding behaviors
4208+
// for any options which are considered to be evaluated from left to right.
4209+
// Specifying any <HostOpts> option which conficts any of the implied options
4210+
// will result in undefined behavior. Potential conflicting options:
4211+
// * Output specification options (-o, -Fo, -Fa, etc)
4212+
// * Phase limiting options (-E, -c, -P, etc)
4213+
void Clang::ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
4214+
const InputInfo &Output,
4215+
const InputInfoList &Inputs,
4216+
const llvm::opt::ArgList &TCArgs) const {
4217+
4218+
// The Host compilation step that occurs here is constructed based on the
4219+
// input from the user. This consists of the compiler to call and the
4220+
// options that will be used during the compilation.
4221+
ArgStringList HostCompileArgs;
4222+
const InputInfo &InputFile = Inputs.front();
4223+
const ToolChain &TC = getToolChain();
4224+
4225+
// Input file.
4226+
HostCompileArgs.push_back(InputFile.getFilename());
4227+
4228+
// When performing the host compilation, we are expecting to only be
4229+
// creating intermediate files, namely preprocessor output, assembly or
4230+
// object files.
4231+
// We are making assumptions in regards to what options are used to
4232+
// generate these intermediate files.
4233+
// gcc/g++/clang/clang++/default | cl
4234+
// Object: -c | -c
4235+
// Preprocessed: -E | -P -Fi<file>
4236+
// Assembly: -S | -c -Fa<file>
4237+
// Header Input: -include <file> | -FI <file>
4238+
//
4239+
// The options used are determined by the compiler name and target triple.
4240+
Arg *HostCompilerDefArg =
4241+
TCArgs.getLastArg(options::OPT_fsycl_host_compiler_EQ);
4242+
assert(HostCompilerDefArg && "Expected host compiler designation.");
4243+
4244+
bool OutputAdded = false;
4245+
StringRef CompilerName =
4246+
llvm::sys::path::stem(HostCompilerDefArg->getValue());
4247+
if (CompilerName.empty())
4248+
TC.getDriver().Diag(diag::err_drv_missing_arg_mtp)
4249+
<< HostCompilerDefArg->getAsString(TCArgs);
4250+
// FIXME: Consider requiring user input to specify a compatibility class
4251+
// to determine the type of host compiler being used.
4252+
SmallVector<StringRef, 4> MSVCCompilers = {"cl", "clang-cl", "icl"};
4253+
bool IsMSVCHostCompiler =
4254+
std::find(MSVCCompilers.begin(), MSVCCompilers.end(), CompilerName) !=
4255+
MSVCCompilers.end();
4256+
4257+
auto addMSVCOutputFile = [&](StringRef Opt) {
4258+
SmallString<128> OutOpt(Opt);
4259+
OutOpt += Output.getFilename();
4260+
HostCompileArgs.push_back(TCArgs.MakeArgString(OutOpt));
4261+
OutputAdded = true;
4262+
};
4263+
// FIXME: Reuse existing toolchains which are already supported to put
4264+
// together the options.
4265+
// FIXME: For any potential obscure host compilers that do not use the
4266+
// 'standard' set of options, we should provide a user interface that allows
4267+
// users to override the implied options.
4268+
if (isa<PreprocessJobAction>(JA)) {
4269+
if (IsMSVCHostCompiler) {
4270+
// Check the output file, if it is 'stdout' we want to use -E.
4271+
if (StringRef(Output.getFilename()).equals("-")) {
4272+
HostCompileArgs.push_back("-E");
4273+
OutputAdded = true;
4274+
} else {
4275+
HostCompileArgs.push_back("-P");
4276+
addMSVCOutputFile("-Fi");
4277+
}
4278+
} else
4279+
HostCompileArgs.push_back("-E");
4280+
} else if (isa<AssembleJobAction>(JA)) {
4281+
HostCompileArgs.push_back("-c");
4282+
if (IsMSVCHostCompiler)
4283+
addMSVCOutputFile("-Fo");
4284+
} else {
4285+
assert((isa<CompileJobAction, BackendJobAction>(JA)) &&
4286+
"Invalid action for external host compilation tool.");
4287+
if (JA.getType() == types::TY_PP_Asm) {
4288+
if (IsMSVCHostCompiler) {
4289+
HostCompileArgs.push_back("-c");
4290+
addMSVCOutputFile("-Fa");
4291+
// The MSVC Compiler does not have a way to just create the assembly
4292+
// file so we create the assembly file and object file, and redirect
4293+
// the object file to a temporary.
4294+
std::string ObjTmpName = C.getDriver().GetTemporaryPath("host", "obj");
4295+
StringRef WrapperFileName =
4296+
C.addTempFile(C.getArgs().MakeArgString(ObjTmpName));
4297+
SmallString<128> ObjOutOpt("-Fo");
4298+
ObjOutOpt += WrapperFileName;
4299+
HostCompileArgs.push_back(C.getArgs().MakeArgString(ObjOutOpt));
4300+
} else
4301+
HostCompileArgs.push_back("-S");
4302+
} else {
4303+
TC.getDriver().Diag(diag::err_drv_output_type_with_host_compiler);
4304+
}
4305+
}
4306+
4307+
// Add default header search directories.
4308+
SmallString<128> BaseDir(C.getDriver().Dir);
4309+
llvm::sys::path::append(BaseDir, "..", "include");
4310+
SmallString<128> SYCLDir(BaseDir);
4311+
llvm::sys::path::append(SYCLDir, "sycl");
4312+
HostCompileArgs.push_back("-I");
4313+
HostCompileArgs.push_back(TCArgs.MakeArgString(SYCLDir));
4314+
HostCompileArgs.push_back("-I");
4315+
HostCompileArgs.push_back(TCArgs.MakeArgString(BaseDir));
4316+
4317+
if (!OutputAdded) {
4318+
// Add output file to the command line. This is assumed to be prefaced
4319+
// with the '-o' option that is used to designate the output file.
4320+
HostCompileArgs.push_back("-o");
4321+
HostCompileArgs.push_back(Output.getFilename());
4322+
}
4323+
4324+
// Add the integration header.
4325+
StringRef Header =
4326+
TC.getDriver().getIntegrationHeader(InputFile.getBaseInput());
4327+
if (types::getPreprocessedType(InputFile.getType()) != types::TY_INVALID &&
4328+
!Header.empty()) {
4329+
HostCompileArgs.push_back(IsMSVCHostCompiler ? "-FI" : "-include");
4330+
HostCompileArgs.push_back(TCArgs.MakeArgString(Header));
4331+
}
4332+
4333+
SmallString<128> ExecPath;
4334+
if (HostCompilerDefArg) {
4335+
ExecPath = HostCompilerDefArg->getValue();
4336+
if (!ExecPath.empty() && ExecPath == llvm::sys::path::stem(ExecPath))
4337+
ExecPath = TC.GetProgramPath(ExecPath.c_str());
4338+
}
4339+
4340+
// Add any user-specified arguments.
4341+
if (Arg *HostCompilerOptsArg =
4342+
TCArgs.getLastArg(options::OPT_fsycl_host_compiler_options_EQ)) {
4343+
SmallVector<const char *, 8> TargetArgs;
4344+
llvm::BumpPtrAllocator BPA;
4345+
llvm::StringSaver S(BPA);
4346+
// Tokenize the string.
4347+
llvm::cl::TokenizeGNUCommandLine(HostCompilerOptsArg->getValue(), S,
4348+
TargetArgs);
4349+
llvm::transform(TargetArgs, std::back_inserter(HostCompileArgs),
4350+
[&TCArgs](StringRef A) { return TCArgs.MakeArgString(A); });
4351+
}
4352+
const Tool *T = TC.SelectTool(JA);
4353+
auto Cmd = std::make_unique<Command>(JA, *T, ResponseFileSupport::None(),
4354+
TCArgs.MakeArgString(ExecPath),
4355+
HostCompileArgs, None);
4356+
4357+
C.addCommand(std::move(Cmd));
4358+
}
4359+
41984360
void Clang::ConstructJob(Compilation &C, const JobAction &JA,
41994361
const InputInfo &Output, const InputInfoList &Inputs,
42004362
const ArgList &Args, const char *LinkingOutput) const {
@@ -4228,6 +4390,14 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
42284390
IsHeaderModulePrecompile || Inputs.size() == 1) &&
42294391
"Unable to handle multiple inputs.");
42304392

4393+
// Perform the SYCL host compilation using an external compiler if the user
4394+
// requested.
4395+
if (Args.hasArg(options::OPT_fsycl_host_compiler_EQ) && IsSYCL &&
4396+
!IsSYCLOffloadDevice) {
4397+
ConstructHostCompilerJob(C, JA, Output, Inputs, Args);
4398+
return;
4399+
}
4400+
42314401
// A header module compilation doesn't have a main input file, so invent a
42324402
// fake one as a placeholder.
42334403
const char *ModuleName = [&]{

‎clang/lib/Driver/ToolChains/Clang.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,6 +88,11 @@ class LLVM_LIBRARY_VISIBILITY Clang : public Tool {
8888
codegenoptions::DebugInfoKind *DebugInfoKind,
8989
bool *EmitCodeView) const;
9090

91+
void ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
92+
const InputInfo &Output,
93+
const InputInfoList &Inputs,
94+
const llvm::opt::ArgList &TCArgs) const;
95+
9196
mutable std::unique_ptr<llvm::raw_fd_ostream> CompilationDatabase = nullptr;
9297
void DumpCompilationDatabase(Compilation &C, StringRef Filename,
9398
StringRef Target,

‎clang/lib/Driver/ToolChains/MSVC.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -435,8 +435,10 @@ void visualstudio::Linker::ConstructJob(Compilation &C, const JobAction &JA,
435435
CmdArgs.push_back("-defaultlib:oldnames");
436436
}
437437

438-
if (!C.getDriver().IsCLMode() && !Args.hasArg(options::OPT_nostdlib) &&
439-
Args.hasArg(options::OPT_fsycl) && !Args.hasArg(options::OPT_nolibsycl)) {
438+
if ((!C.getDriver().IsCLMode() && !Args.hasArg(options::OPT_nostdlib) &&
439+
Args.hasArg(options::OPT_fsycl) &&
440+
!Args.hasArg(options::OPT_nolibsycl)) ||
441+
Args.hasArg(options::OPT_fsycl_host_compiler_EQ)) {
440442
if (Args.hasArg(options::OPT__SLASH_MDd))
441443
CmdArgs.push_back("-defaultlib:sycld.lib");
442444
else

‎clang/test/Driver/sycl-host-compiler.cpp

Lines changed: 64 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,64 @@
1+
// Tests the abilities involved with using an external host compiler
2+
// REQUIRES: clang-driver
3+
4+
/// enabling with -fsycl-host-compiler
5+
// RUN: %clangxx -fsycl -fsycl-host-compiler=/some/dir/g++ %s -### 2>&1 \
6+
// RUN: | FileCheck -check-prefix=HOST_COMPILER %s
7+
// HOST_COMPILER: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]"
8+
// HOST_COMPILER: g++{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-o" "[[HOSTOBJ:.+\.o]]"{{.*}} "-include" "[[INTHEADER]]"
9+
// HOST_COMPILER: ld{{.*}} "[[HOSTOBJ]]"
10+
11+
// RUN: %clang_cl -fsycl -fsycl-host-compiler=/some/dir/cl %s -### 2>&1 \
12+
// RUN: | FileCheck -check-prefix=HOST_COMPILER_CL %s
13+
// HOST_COMPILER_CL: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]"
14+
// HOST_COMPILER_CL: cl{{.*}} "-Fo[[HOSTOBJ:.+\.obj]]"{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-FI" "[[INTHEADER]]"
15+
// HOST_COMPILER_CL: link{{.*}} "[[HOSTOBJ]]"
16+
17+
/// check for additional host options
18+
// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -fsycl-host-compiler-options="-DFOO -DBAR" %s -### 2>&1 \
19+
// RUN: | FileCheck -check-prefix=HOST_OPTIONS %s
20+
// HOST_OPTIONS: g++{{.*}} "-o" "[[HOSTOBJ:.+\.o]]"{{.*}} "-DFOO" "-DBAR"
21+
22+
// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -fsycl-host-compiler-options="/DFOO /DBAR /O2" %s -### 2>&1 \
23+
// RUN: | FileCheck -check-prefix=HOST_OPTIONS_CL %s
24+
// HOST_OPTIONS_CL: cl{{.*}} "-Fo[[HOSTOBJ:.+\.obj]]"{{.*}} "/DFOO" "/DBAR" "/O2"
25+
26+
/// preprocessing
27+
// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -E %s -### 2>&1 \
28+
// RUN: | FileCheck -check-prefix=HOST_PREPROCESS %s
29+
// HOST_PREPROCESS: g++{{.*}} "-E"{{.*}} "-o" "[[PPOUT:.+\.ii]]"
30+
// HOST_PREPROCESS: g++{{.*}} "-E"{{.*}} "-o" "-"
31+
// HOST_PREPROCESS: clang-offload-bundler{{.*}} "-inputs={{.*}}.ii,[[PPOUT]]"
32+
33+
// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -E %s -### 2>&1 \
34+
// RUN: | FileCheck -check-prefix=HOST_PREPROCESS_CL %s
35+
// HOST_PREPROCESS_CL: cl{{.*}} "-P"{{.*}} "-Fi[[PPOUT:.+\.ii]]"
36+
// HOST_PREPROCESS_CL: cl{{.*}} "-E"
37+
// HOST_PREPROCESS_CL: clang-offload-bundler{{.*}} "-inputs={{.*}}.ii,[[PPOUT]]"
38+
39+
/// obj output
40+
// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -c %s -### 2>&1 \
41+
// RUN: | FileCheck -check-prefix=HOST_OBJECT %s
42+
// HOST_OBJECT: g++{{.*}} "-c"{{.*}} "-o" "[[OBJOUT:.+\.o]]"
43+
// HOST_OBJECT: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[OBJOUT]]"
44+
45+
// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -c %s -### 2>&1 \
46+
// RUN: | FileCheck -check-prefix=HOST_OBJECT_CL %s
47+
// HOST_OBJECT_CL: cl{{.*}} "-c"{{.*}} "-Fo[[OBJOUT:.+\.obj]]"
48+
// HOST_OBJECT_CL: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[OBJOUT]]"
49+
50+
/// assembly output
51+
// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -S %s -### 2>&1 \
52+
// RUN: | FileCheck -check-prefix=HOST_ASSEMBLY %s
53+
// HOST_ASSEMBLY: g++{{.*}} "-S"{{.*}} "-o" "[[ASMOUT:.+\.s]]"
54+
// HOST_ASSEMBLY: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[ASMOUT]]"
55+
56+
// RUN: %clangxx -fsycl -fsycl-host-compiler=cl -S %s -### 2>&1 \
57+
// RUN: | FileCheck -check-prefix=HOST_ASSEMBLY_CL %s
58+
// HOST_ASSEMBLY_CL: cl{{.*}} "-c"{{.*}} "-Fa[[ASMOUT:.+\.s]]" "-Fo{{.*}}.obj"
59+
// HOST_ASSEMBLY_CL: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[ASMOUT]]"
60+
61+
/// missing argument error -fsycl-host-compiler=
62+
// RUN: %clangxx -fsycl -fsycl-host-compiler= -c -### %s 2>&1 \
63+
// RUN: | FileCheck -check-prefix=HOST_COMPILER_NOARG %s
64+
// HOST_COMPILER_NOARG: missing argument to '-fsycl-host-compiler='

‎sycl/doc/UsersManual.md

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -222,6 +222,26 @@ and not recommended to use in production environment.
222222
same as specifying -fsycl-help with no argument and emits help for all
223223
backends.
224224

225+
**`-fsycl-host-compiler=<arg>`**
226+
227+
Informs the compiler driver that the host compilation step that is performed
228+
as part of the greater compilation flow will be performed by the compiler
229+
<arg>. It is expected that <arg> is the compiler to be called, either by
230+
name (in which the PATH will be used to discover it) or a fully qualified
231+
directory with compiler to invoke. This option is only useful when -fsycl
232+
is provided on the command line.
233+
234+
**`-fsycl-host-compiler-options="opts"`**
235+
236+
Passes along the space separated quoted "opts" string as option arguments
237+
to the compiler specified with the -fsycl-host-compiler=<arg> option. It is
238+
expected that the options used here are compatible with the compiler
239+
specified via -fsycl-host-compiler=<arg>.
240+
241+
NOTE: Using -fsycl-host-compiler-options to pass any kind of phase limiting
242+
options (e.g. -c, -E, -S) may interfere with the expected output set during
243+
the host compilation. Doing so is considered undefined behavior.
244+
225245
# Example: SYCL device code compilation
226246

227247
To invoke SYCL device compiler set `-fsycl-device-only` flag.

‎sycl/test/regression/fsycl-host-compiler-win.cpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -DDEFINE_CHECK -fsycl-host-compiler-options="-DDEFINE_CHECK" /Fe%t1.exe %s
2+
// RUN: %RUN_ON_HOST %t1.exe
3+
// REQUIRES: system-windows
4+
//
5+
//==------- fsycl-host-compiler-win.cpp - external host compiler test ------==//
6+
//
7+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
8+
// See https://llvm.org/LICENSE.txt for license information.
9+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
10+
//
11+
//===----------------------------------------------------------------------===//
12+
//
13+
// Uses -fsycl-host-compiler=<compiler> on a simple test, requires 'cl'
14+
15+
#include <CL/sycl.hpp>
16+
17+
#ifndef DEFINE_CHECK
18+
#error predefined macro not set
19+
#endif // DEFINE_CHECK
20+
21+
using namespace cl::sycl;
22+
23+
int main() {
24+
int data[] = {0, 0, 0};
25+
26+
{
27+
buffer<int, 1> b(data, range<1>(3), {property::buffer::use_host_ptr()});
28+
queue q;
29+
q.submit([&](handler &cgh) {
30+
auto B = b.get_access<access::mode::write>(cgh);
31+
cgh.parallel_for<class test>(range<1>(3), [=](id<1> idx) {
32+
B[idx] = 1;
33+
});
34+
});
35+
}
36+
37+
bool isSuccess = true;
38+
39+
for (int i = 0; i < 3; i++)
40+
if (data[i] != 1) isSuccess = false;
41+
42+
if (!isSuccess)
43+
return -1;
44+
45+
return 0;
46+
}

0 commit comments

Comments
 (0)