Issue 1:
System information

Have I written custom code (as opposed to using a stock example script provided in TensorFlow): N/A

OS Platform and Distribution (e.g., Linux Ubuntu 16.04): macOS 'High Sierra' Version 10.13.4 (17E202)

TensorFlow installed from (source or binary): source

TensorFlow version (use command below): N/A, attempting to compile at e365dea with minimal local changes (see below)

Python version: 2.7.15

Bazel version (if compiling from source): 0.13.1-homebrew

GCC/Compiler version (if compiling from source): Apple LLVM version 8.1.0 (clang-802.0.42)

CUDA/cuDNN version: cuda_9.2.64_mac with cuda_9.2.64.1_mac, cudnn-9.2-osx-x64-v7.1

GPU model and memory: NVIDIA GeForce GT 750M with 2 GB device memory, CUDA Compute Capability 3.0

Exact command to reproduce:

./configure
bazel build --config=opt --config=cuda --save_temps --explain=explain.txt --verbose_explanations --verbose_failures --linkopt=-Wl,-rpath,/usr/local/cuda/lib //tensorflow/tools/pip_package:build_pip_package

Describe the problem

This is a re-occurrence of #9072, except that the solutions mentioned there (not using clang as the CUDA compiler, using CommandLineTools) do not resolve the problem.

To configure, I selected the following:

You have bazel 0.13.1-homebrew installed.
Please specify the location of python. [Default is /usr/local/opt/python@2/bin/python2.7]:

Found possible Python library paths:
/usr/local/Cellar/python@2/2.7.15/Frameworks/Python.framework/Versions/2.7/lib/python2.7/site-packages
Please input the desired Python library path to use. Default is [/usr/local/Cellar/python@2/2.7.15/Frameworks/Python.framework/Versions/2.7/lib/python2.7/site-packages]

Do you wish to build TensorFlow with Google Cloud Platform support? [Y/n]:
Google Cloud Platform support will be enabled for TensorFlow.

Do you wish to build TensorFlow with Hadoop File System support? [Y/n]:
Hadoop File System support will be enabled for TensorFlow.

Do you wish to build TensorFlow with Amazon S3 File System support? [Y/n]:
Amazon S3 File System support will be enabled for TensorFlow.

Do you wish to build TensorFlow with Apache Kafka Platform support? [Y/n]:
Apache Kafka Platform support will be enabled for TensorFlow.

Do you wish to build TensorFlow with XLA JIT support? [y/N]:
No XLA JIT support will be enabled for TensorFlow.

Do you wish to build TensorFlow with GDR support? [y/N]:
No GDR support will be enabled for TensorFlow.

Do you wish to build TensorFlow with VERBS support? [y/N]:
No VERBS support will be enabled for TensorFlow.

Do you wish to build TensorFlow with OpenCL SYCL support? [y/N]:
No OpenCL SYCL support will be enabled for TensorFlow.

Do you wish to build TensorFlow with CUDA support? [y/N]: y
CUDA support will be enabled for TensorFlow.

Please specify the CUDA SDK version you want to use. [Leave empty to default to CUDA 9.0]: 9.2

Please specify the location where CUDA 9.2 toolkit is installed. Refer to README.md for more details. [Default is /usr/local/cuda]:

Please specify the cuDNN version you want to use. [Leave empty to default to cuDNN 7.0]: 7.1.4

Please specify the location where cuDNN 7 library is installed. Refer to README.md for more details. [Default is /usr/local/cuda]:

Please specify a list of comma-separated Cuda compute capabilities you want to build with.
You can find the compute capability of your device at: https://developer.nvidia.com/cuda-gpus.
Please note that each additional compute capability significantly increases your build time and binary size. [Default is: 3.5,5.2]3.0

Do you want to use clang as CUDA compiler? [y/N]:
nvcc will be used as CUDA compiler.

Please specify which gcc should be used by nvcc as the host compiler. [Default is /usr/bin/gcc]:

Do you wish to build TensorFlow with MPI support? [y/N]:
No MPI support will be enabled for TensorFlow.

Please specify optimization flags to use during compilation when bazel option "--config=opt" is specified [Default is -march=native]:

Would you like to interactively configure ./WORKSPACE for Android builds? [y/N]:
Not configuring the WORKSPACE for Android builds.

Preconfigured Bazel build configs. You can use any of the below by adding "--config=<>" to your build command. See tools/bazel.rc for more details.
--config=mkl	# Build with MKL support.
--config=monolithic	# Config for mostly static monolithic build.
Configuration finished

To build, I try:

bazel build --config=opt --config=cuda --save_temps --explain=explain.txt --verbose_explanations --verbose_failures --linkopt=-Wl,-rpath,/usr/local/cuda/lib //tensorflow/tools/pip_package:build_pip_package

However, this results in the error:

Starting local Bazel server and connecting to it...
............
WARNING: The following configs were expanded more than once: [cuda]. For repeatable flags, repeats are counted twice and may lead to unexpected behavior.
ERROR: Inconsistent crosstool configuration; no toolchain corresponding to 'local_darwin' found for cpu 'darwin'
INFO: Elapsed time: 0.903s
INFO: 0 processes.
FAILED: Build did NOT complete successfully (2 packages loaded)

The output of xcode-select -p is:

/Library/Developer/CommandLineTools

The output of /usr/bin/gcc --version is:

Configured with: --prefix=/Library/Developer/CommandLineTools/usr --with-gxx-include-dir=/usr/include/c++/4.2.1
Apple LLVM version 8.1.0 (clang-802.0.42)
Target: x86_64-apple-darwin17.5.0
Thread model: posix
InstalledDir: /Library/Developer/CommandLineTools/usr/bin

I am able to build & run the 'deviceQuery' CUDA SDK sample without issue.

Source code / logs

The only local changes from e365dea I have are:

diff --git a/tensorflow/core/kernels/concat_lib_gpu_impl.cu.cc b/tensorflow/core/kernels/concat_lib_gpu_impl.cu.cc
index a561d918bd..46c91b4511 100644
--- a/tensorflow/core/kernels/concat_lib_gpu_impl.cu.cc
+++ b/tensorflow/core/kernels/concat_lib_gpu_impl.cu.cc
@@ -69,7 +69,7 @@ __global__ void concat_variable_kernel(
IntType num_inputs = input_ptr_data.size;

// verbose declaration needed due to template
- extern __shared__ __align__(sizeof(T)) unsigned char smem[];
+ extern __shared__ __align__(sizeof(T) > 16 ? sizeof(T) : 16) unsigned char smem[];
IntType* smem_col_scan = reinterpret_cast<IntType*>(smem);

if (useSmem) {
diff --git a/tensorflow/core/kernels/depthwise_conv_op_gpu.cu.cc b/tensorflow/core/kernels/depthwise_conv_op_gpu.cu.cc
index 5390222b3a..fcbd733614 100644
--- a/tensorflow/core/kernels/depthwise_conv_op_gpu.cu.cc
+++ b/tensorflow/core/kernels/depthwise_conv_op_gpu.cu.cc
@@ -172,7 +172,7 @@ __global__ __launch_bounds__(1024, 2) void DepthwiseConv2dGPUKernelNHWCSmall(
const DepthwiseArgs args, const T* input, const T* filter, T* output) {
assert(CanLaunchDepthwiseConv2dGPUSmall(args));
// Holds block plus halo and filter data for blockDim.x depths.
- extern __shared__ __align__(sizeof(T)) unsigned char shared_memory[];
+ extern __shared__ __align__(sizeof(T) > 16 ? sizeof(T) : 16) unsigned char shared_memory[];
T* const shared_data = reinterpret_cast<T*>(shared_memory);

const int num_batches = args.batch;
@@ -452,7 +452,7 @@ __global__ __launch_bounds__(1024, 2) void DepthwiseConv2dGPUKernelNCHWSmall(
const DepthwiseArgs args, const T* input, const T* filter, T* output) {
assert(CanLaunchDepthwiseConv2dGPUSmall(args));
// Holds block plus halo and filter data for blockDim.z depths.
- extern __shared__ __align__(sizeof(T)) unsigned char shared_memory[];
+ extern __shared__ __align__(sizeof(T) > 16 ? sizeof(T) : 16) unsigned char shared_memory[];
T* const shared_data = reinterpret_cast<T*>(shared_memory);

const int num_batches = args.batch;
@@ -1118,7 +1118,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNHWCSmall(
const DepthwiseArgs args, const T* output, const T* input, T* filter) {
assert(CanLaunchDepthwiseConv2dBackpropFilterGPUSmall(args, blockDim.z));
// Holds block plus halo and filter data for blockDim.x depths.
- extern __shared__ __align__(sizeof(T)) unsigned char shared_memory[];
+ extern __shared__ __align__(sizeof(T) > 16 ? sizeof(T) : 16) unsigned char shared_memory[];
T* const shared_data = reinterpret_cast<T*>(shared_memory);

const int num_batches = args.batch;
@@ -1388,7 +1388,7 @@ __launch_bounds__(1024, 2) void DepthwiseConv2dBackpropFilterGPUKernelNCHWSmall(
const DepthwiseArgs args, const T* output, const T* input, T* filter) {
assert(CanLaunchDepthwiseConv2dBackpropFilterGPUSmall(args, blockDim.x));
// Holds block plus halo and filter data for blockDim.z depths.
- extern __shared__ __align__(sizeof(T)) unsigned char shared_memory[];
+ extern __shared__ __align__(sizeof(T) > 16 ? sizeof(T) : 16) unsigned char shared_memory[];
T* const shared_data = reinterpret_cast<T*>(shared_memory);

const int num_batches = args.batch;
diff --git a/tensorflow/core/kernels/split_lib_gpu.cu.cc b/tensorflow/core/kernels/split_lib_gpu.cu.cc
index 393818730b..58a1294005 100644
--- a/tensorflow/core/kernels/split_lib_gpu.cu.cc
+++ b/tensorflow/core/kernels/split_lib_gpu.cu.cc
@@ -121,7 +121,7 @@ __global__ void split_v_kernel(const T* input_ptr,
int num_outputs = output_ptr_data.size;

// verbose declaration needed due to template
- extern __shared__ __align__(sizeof(T)) unsigned char smem[];
+ extern __shared__ __align__(sizeof(T) > 16 ? sizeof(T) : 16) unsigned char smem[];
IntType* smem_col_scan = reinterpret_cast<IntType*>(smem);

if (useSmem) {
diff --git a/third_party/gpus/cuda/BUILD.tpl b/third_party/gpus/cuda/BUILD.tpl
index 2a37c65bc7..43446dd99b 100644
--- a/third_party/gpus/cuda/BUILD.tpl
+++ b/third_party/gpus/cuda/BUILD.tpl
@@ -110,7 +110,7 @@ cc_library(
".",
"cuda/include",
],
- linkopts = ["-lgomp"],
+ #linkopts = ["-lgomp"],
linkstatic = 1,
visibility = ["//visibility:public"],
)
diff --git a/third_party/toolchains/gpus/cuda/BUILD b/third_party/toolchains/gpus/cuda/BUILD
index 4cb8380938..d025c4f3aa 100644
--- a/third_party/toolchains/gpus/cuda/BUILD
+++ b/third_party/toolchains/gpus/cuda/BUILD
@@ -115,7 +115,7 @@ cc_library(
".",
"cuda/include",
],
- linkopts = ["-lgomp"],
+ #linkopts = ["-lgomp"],
linkstatic = 1,
visibility = ["//visibility:public"],
)

Issue 2:
I encountered a problem when using a self-trained face-recognition model to make inference on android platform (using c++ api, just like the android demo). The error says something like this:

06-05 16:25:11.322 28605-28605/jp.narr.tensorflowmnist I/native: tensorflow_jni.cc:196 End computing.
06-05 16:25:11.322 28605-28605/jp.narr.tensorflowmnist E/native: tensorflow_jni.cc:199 Error during inference: Invalid argument: No OpKernel was registered to support Op 'Inv' with these attrs
[[Node: incept5b/in4_conv1x1_55/batch_norm/moments/moments_1/divisor = Inv[T=DT_FLOAT](incept5b/in4_conv1x1_55/batch_norm/moments/moments/Const)]]
06-05 16:25:11.322 28605-28605/jp.narr.tensorflowmnist A/libc: Fatal signal 11 (SIGSEGV), code 1, fault addr 0x10 in tid 28605 (tensorflowmnist)
06-05 16:25:11.423 186-186/? I/DEBUG: *** *** *** *** *** *** *** *** *** *** *** ***

It is similar to the issue #1269

I don't understand why it causes an error?

All the other layers ( from incept3a to incept5a) have almost the same structures, but there's no error....

Could anyone give me some advice?

Thanks a lot!

The structure of the model I use is like this:

def inference_nn4_max_pool_96(images, pool_type, use_lrn, keep_probability, phase_train=True):
conv1 = _conv(images, 3, 64, 7, 7, 2, 2, 'SAME', 'conv1_7x7', phase_train=phase_train, use_batch_norm=True)
pool1 = _mpool(conv1, 3, 3, 2, 2, 'SAME')
if use_lrn:
lrn1 = tf.nn.local_response_normalization(pool1, depth_radius=5, bias=1.0, alpha=1e-4, beta=0.75)
else:
lrn1 = pool1
conv2 = _conv(lrn1, 64, 64, 1, 1, 1, 1, 'SAME', 'conv2_1x1', phase_train=phase_train, use_batch_norm=True)
conv3 = _conv(conv2, 64, 192, 3, 3, 1, 1, 'SAME', 'conv3_3x3', phase_train=phase_train, use_batch_norm=True)
if use_lrn:
lrn2 = tf.nn.local_response_normalization(conv3, depth_radius=5, bias=1.0, alpha=1e-4, beta=0.75)
else:
lrn2 = conv3
pool3 = _mpool(lrn2, 3, 3, 2, 2, 'SAME')

incept3a = _inception(pool3, 192, 1, 64, 96, 128, 16, 32, 3, 32, 1, 'MAX', 'incept3a', phase_train=phase_train, use_batch_norm=True)
incept3b = _inception(incept3a, 256, 1, 64, 96, 128, 32, 64, 3, 64, 1, pool_type, 'incept3b', phase_train=phase_train, use_batch_norm=True)
incept3c = _inception(incept3b, 320, 2, 0, 128, 256, 32, 64, 3, 0, 2, 'MAX', 'incept3c', phase_train=phase_train, use_batch_norm=True)

incept4a = _inception(incept3c, 640, 1, 256, 96, 192, 32, 64, 3, 128, 1, pool_type, 'incept4a', phase_train=phase_train, use_batch_norm=True)
incept4b = _inception(incept4a, 640, 1, 224, 112, 224, 32, 64, 3, 128, 1, pool_type, 'incept4b', phase_train=phase_train, use_batch_norm=True)
incept4c = _inception(incept4b, 640, 1, 192, 128, 256, 32, 64, 3, 128, 1, pool_type, 'incept4c', phase_train=phase_train, use_batch_norm=True)
incept4d = _inception(incept4c, 640, 1, 160, 144, 288, 32, 64, 3, 128, 1, pool_type, 'incept4d', phase_train=phase_train, use_batch_norm=True)
incept4e = _inception(incept4d, 640, 2, 0, 160, 256, 64, 128, 3, 0, 2, 'MAX', 'incept4e', phase_train=phase_train, use_batch_norm=True)

incept5a = _inception(incept4e, 1024, 1, 384, 192, 384, 0, 0, 3, 128, 1, pool_type, 'incept5a', phase_train=phase_train, use_batch_norm=True)
incept5b = _inception(incept5a, 896, 1, 384, 192, 384, 0, 0, 3, 128, 1, 'MAX', 'incept5b', phase_train=phase_train, use_batch_norm=True)
pool6 = _apool(incept5b, 3, 3, 1, 1, 'VALID')

resh1 = tf.reshape(pool6, [-1, 896])
affn1 = _affine(resh1, 896, 128)
if keep_probability<1.0:
affn1 = control_flow_ops.cond(phase_train,
lambda: tf.nn.dropout(affn1, keep_probability), lambda: affn1)
norm = tf.nn.l2_normalize(affn1, 1, 1e-10, name='embeddings')

return norm