This is an archive of the discontinued LLVM Phabricator instance.

[libomptarget] Replace forceinline with inline to work around nvcc deleting forceinline symbols
AbandonedPublic

Authored by JonChesterfield on Nov 5 2019, 10:42 AM.

Download Raw Diff

Details

Reviewers

ABataev
jdoerfert
grokos
ikitayama
tianshilei1992

Summary

[libomptarget] Replace forceinline with inline to work around nvcc deleting forceinline symbols

There is a problem with some installs of nvcc deleting the forceinline symbols from support.cu.
Adding the people who reported this as reviewers. This patch is an eductated guess at a fix,
posting it without a local repro for the failure. See D69857 for an alternative.

Diff Detail

Repository

rG LLVM Github Monorepo

Build Status

Buildable 40530
Build 40642: arc lint + arc unit

Event Timeline

JonChesterfield created this revision.Nov 5 2019, 10:42 AM

Herald added a project: Restricted Project. · View Herald TranscriptNov 5 2019, 10:42 AM

Herald added a subscriber: openmp-commits. · View Herald Transcript

JonChesterfield mentioned this in D69857: [libomptarget] Replace INLINE with DEVICE to work around nvcc deleting forceinline symbols.Nov 5 2019, 10:42 AM

Harbormaster completed remote builds in B40530: Diff 227921.Nov 5 2019, 10:47 AM

Nope, it does not work, still unresolved symbols.

In D69859#1734445, @ABataev wrote:

Nope, it does not work, still unresolved symbols.

That's a shame. How about D69857?

JonChesterfield mentioned this in D69885: [libomptarget] Revert all improvements to support.Nov 6 2019, 2:51 AM

JonChesterfield mentioned this in rG7cea0cea77d1: [libomptarget] Revert all improvements to support.Nov 6 2019, 7:50 AM

JonChesterfield abandoned this revision.Nov 12 2019, 8:36 AM

Revision Contents

Path

Size

openmp/

libomptarget/

deviceRTLs/

nvptx/

src/

target_impl.h

2 lines

Diff 227921

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

	Show All 10 Lines
	//===----------------------------------------------------------------------===//			//===----------------------------------------------------------------------===//
	#ifndef _TARGET_IMPL_H_			#ifndef _TARGET_IMPL_H_
	#define _TARGET_IMPL_H_			#define _TARGET_IMPL_H_

	#include <cuda.h>			#include <cuda.h>
	#include "nvptx_interface.h"			#include "nvptx_interface.h"

	#define DEVICE __device__			#define DEVICE __device__
	#define INLINE __forceinline__ DEVICE			#define INLINE __inline__ DEVICE
	#define NOINLINE __noinline__ DEVICE			#define NOINLINE __noinline__ DEVICE

	////////////////////////////////////////////////////////////////////////////////			////////////////////////////////////////////////////////////////////////////////
	// Kernel options			// Kernel options
	////////////////////////////////////////////////////////////////////////////////			////////////////////////////////////////////////////////////////////////////////

	////////////////////////////////////////////////////////////////////////////////			////////////////////////////////////////////////////////////////////////////////
	// The following def must match the absolute limit hardwired in the host RTL			// The following def must match the absolute limit hardwired in the host RTL
	▲ Show 20 Lines • Show All 137 Lines • Show Last 20 Lines