diff --git a/CHANGELOG.md b/CHANGELOG.md new file mode 100644 --- /dev/null +++ b/CHANGELOG.md @@ -0,0 +1,5 @@ +# Revision history for nvvm + +## 0.7.5.0 -- YYYY-mm-dd + +* First version. Released on an unsuspecting world. diff --git a/Foreign/NVVM.hs b/Foreign/NVVM.hs new file mode 100644 --- /dev/null +++ b/Foreign/NVVM.hs @@ -0,0 +1,190 @@ +-------------------------------------------------------------------------------- +-- | +-- Module : Foreign.NVVM +-- Copyright : [2016] Trevor L. McDonell +-- License : BSD +-- +-- This module defines an interface to the /libNVVM/ library provided by NVIDIA as +-- part of the CUDA toolkit. It compiles NVVM IR, a compiler intermediate +-- representation based on LLVM IR, into PTX code suitable for execution on +-- NVIDIA GPUs. NVVM IR is a subset of LLVM IR, with a set of rules, +-- restrictions, conventions, and intrinsic functions. +-- +-- NVIDIA's own 'nvcc' compiler uses NVVM IR and /libNVVM/ internally as part of +-- the CUDA C compilation process. In contrast to the (open-source) NVPTX target +-- included with the standard LLVM toolchain (which also compiles NVVM IR into +-- PTX code), /libNVVM/ includes a set of proprietary optimisation passes, which +-- /may/ result in faster GPU code. More information on NVVM IR can be found +-- here: +-- +-- +-- +-- The following is a short tutorial on using this library. The steps can be +-- copied into a file, or run directly in @ghci@, in which case @ghci@ should be +-- launched with the option @-fno-ghci-sandbox@. This is because CUDA maintains +-- CPU-local state, so operations must be run from a bound thread. +-- +-- Note that the focus of this library is the generation of executable PTX code +-- from NVVM IR, so we will additionally need to use the 'cuda' package to +-- control and execute the compiled program. In this tutorial we will go over +-- those steps quickly, but see the 'cuda' package for more information and +-- a similar tutorial: +-- +-- +-- +-- +-- [/Initialise the CUDA environment/] +-- +-- Before any operation can be performed, we must initialise the CUDA Driver +-- API. +-- +-- >>> import Foreign.CUDA.Driver as CUDA +-- >>> CUDA.initialise [] +-- +-- Select a GPU and create an execution context for that device. Each available +-- device is given a unique numeric identifier (beginning at zero). For this +-- example we just select the first device (the default). +-- +-- >>> dev0 <- CUDA.device 0 +-- >>> prp0 <- CUDA.props dev0 +-- >>> ctx0 <- CUDA.create dev0 [] +-- +-- Remember that once the context is no longer needed, it should be +-- 'Foreign.CUDA.Driver.Context.Base.destroy'ed in order to free up any +-- resources that were allocated into it. +-- +-- [/Compiling kernels with NVVM/] +-- +-- For this example we will step through executing the equivalent of the +-- following Haskell function, which element-wise adds the elements of two +-- arrays: +-- +-- >>> vecAdd xs ys = zipWith (+) xs ys +-- +-- The following NVVM IR implements this for the GPU. Note that this example is +-- written using NVVM IR version 1.2 syntax (corresponding to CUDA toolkit 7.5), +-- which is based on LLVM IR version 3.4. The human readable representation of +-- LLVM IR (and by extension NVVM IR) is notorious for changing between +-- releases, whereas the bitcode representation is somewhat more stable. You may +-- wish to keep this in mind for your own programs. +-- +-- > target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" +-- > target triple = "nvptx64-nvidia-cuda" +-- > +-- > define void @vecAdd(float* %A, float* %B, float* %C) { +-- > entry: +-- > ; What is my ID? +-- > %id = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind +-- > +-- > ; Compute pointers into A, B, and C +-- > %ptrA = getelementptr float* %A, i32 %id +-- > %ptrB = getelementptr float* %B, i32 %id +-- > %ptrC = getelementptr float* %C, i32 %id +-- > +-- > ; Read A, B +-- > %valA = load float* %ptrA, align 4 +-- > %valB = load float* %ptrB, align 4 +-- > +-- > ; Compute C = A + B +-- > %valC = fadd float %valA, %valB +-- > +-- > ; Store back to C +-- > store float %valC, float* %ptrC, align 4 +-- > +-- > ret void +-- > } +-- > +-- > ; Intrinsic to read threadIdx.x +-- > declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() readnone nounwind +-- > +-- > !nvvm.annotations = !{!0} +-- > !0 = metadata !{void (float*, float*, float*)* @vecAdd, metadata !"kernel", i64 1} +-- +-- For reference, in CUDA this kernel would have been written as: +-- +-- > extern "C" __global__ void vecAdd(float *xs, float* ys, float *zs) +-- > { +-- > int ix = threadIdx.x; +-- > +-- > zs[ix] = xs[ix] + ys[ix]; +-- > } +-- +-- The NVVM IR can be stored directly in the program as a @[Byte]String@, but +-- here I will assume that it is saved to a file @vector_add.ll@: +-- +-- >>> import Data.ByteString as B +-- >>> ll <- B.readFile "vector_add.ll" +-- +-- Now we can use NVVM to compile this into PTX code: +-- +-- >>> import Foreign.NVVM as NVVM +-- >>> ptx <- NVVM.compileModule "vecAdd" ll [ NVVM.Target (CUDA.computeCapability prp0) ] +-- +-- Notice that we asked NVVM to specialise the generated PTX code for our +-- current device. By default the code will be compiled for compute capability +-- 2.0 (the earliest supported target). +-- +-- The generated PTX code can then be loaded into the current CUDA execution +-- context, from which we can extract a reference to the GPU kernel that we will +-- later execute. +-- +-- >>> mdl <- CUDA.loadData (NVVM.compileResult ptx) +-- >>> vecAdd <- CUDA.getFun mdl "vecAdd" +-- +-- After we are finished with the module, it is a good idea to +-- 'Foreign.CUDA.Driver.Module.Base.unload' it in order to free any resources it +-- used. +-- +-- [/Executing the kernel/] +-- +-- Executing the 'vecAdd' kernel now proceeds exactly like executing any other +-- kernel function using the CUDA Driver API. The following is a quick overview; +-- see the tutorial in the 'cuda' package for more information. +-- +-- First, generate some data and copy it to the device. We also allocate an +-- (uninitialised) array on the device to store the results. +-- +-- >>> let xs = [1..256] :: [Float] +-- >>> let ys = [2,4..512] :: [Float] +-- >>> xs_dev <- CUDA.newListArray xs +-- >>> ys_dev <- CUDA.newListArray ys +-- >>> zs_dev <- CUDA.mallocArray 256 :: IO (CUDA.DevicePtr Float) +-- +-- For this simple kernel we execute it using a single (one dimensional) thread +-- block, with one thread computing each element of the output. +-- +-- >>> CUDA.launchKernel vecAdd (1,1,1) (256,1,1) 0 Nothing [CUDA.VArg xs_dev, CUDA.VArg ys_dev, CUDA.VArg zs_dev] +-- +-- Finally, we can copy the results back to the host, and deallocate the arrays +-- from the GPU. +-- +-- >>> zs <- CUDA.peekListArray 256 zs_dev +-- >>> CUDA.free xs_dev +-- >>> CUDA.free ys_dev +-- >>> CUDA.free zs_dev +-- +-- +-- [/Next steps/] +-- +-- The library also provides functions for compiling several NVVM IR sources +-- into a single module. In particular this is useful when linking against +-- @libdevice@, a standard library of functions in NVVM IR which implement, for +-- example, math primitives and bitwise operations. More information on +-- @libdevice@ can be found here: +-- +-- +-- +-------------------------------------------------------------------------------- + +module Foreign.NVVM ( + + module Foreign.NVVM.Compile, + module Foreign.NVVM.Error, + module Foreign.NVVM.Info, + +) where + +import Foreign.NVVM.Compile +import Foreign.NVVM.Error +import Foreign.NVVM.Info + diff --git a/Foreign/NVVM/Compile.chs b/Foreign/NVVM/Compile.chs new file mode 100644 --- /dev/null +++ b/Foreign/NVVM/Compile.chs @@ -0,0 +1,283 @@ +{-# LANGUAGE BangPatterns #-} +{-# OPTIONS_GHC -funbox-strict-fields #-} +-------------------------------------------------------------------------------- +-- | +-- Module : Foreign.NVVM.Compile +-- Copyright : [2016] Trevor L. McDonell +-- License : BSD +-- +-- Program compilation +-- +-------------------------------------------------------------------------------- + +module Foreign.NVVM.Compile ( + + Program, + Result(..), + CompileOption(..), + + compileModule, compileModules, + + create, + destroy, + addModule, addModuleFromPtr, + compile, + verify + +) where + +import Foreign.CUDA.Analysis +import Foreign.NVVM.Error +import Foreign.NVVM.Internal.C2HS + +import Foreign.C +import Foreign.Marshal +import Foreign.Ptr +import Foreign.ForeignPtr +import Foreign.Storable + +import Control.Exception +import Data.Word +import Data.ByteString ( ByteString ) +import Text.Printf +import qualified Data.ByteString.Char8 as B +import qualified Data.ByteString.Unsafe as B +import qualified Data.ByteString.Internal as B + + +#include "cbits/stubs.h" +{# context lib="nvvm" #} + + +-- | An NVVM program +-- +newtype Program = Program { useProgram :: {# type nvvmProgram #} } + deriving ( Eq, Show ) + +-- | The result of compiling an NVVM program. +-- +data Result = Result + { compileResult :: !ByteString -- ^ The compiled kernel, which can be loaded into the current program using 'Foreign.CUDA.Driver.loadData*' + , compileLog :: !ByteString -- ^ Warning messages generated by the compiler/verifier + } + +-- | Program compilation options +-- +data CompileOption + = OptimisationLevel !Int -- ^ optimisation level, from 0 (disable optimisations) to 3 (default) + | Target !Compute -- ^ target architecture to compile for (default: compute 2.0) + | FlushToZero -- ^ flush denormal values to zero when performing single-precision floating-point operations (default: no) + | NoFMA -- ^ disable fused-multiply-add instructions (default: enabled) + | FastSqrt -- ^ use a fast approximation for single-precision floating-point square root (default: no) + | FastDiv -- ^ use a fast approximation for single-precision floating-point division and reciprocal (default: no) + | GenerateDebugInfo -- ^ generate debugging information (-g) (default: no) + deriving ( Eq, Show ) + + +-- | Compile an NVVM IR module, in either bitcode or textual representation, +-- into PTX code. +-- +{-# INLINEABLE compileModule #-} +compileModule + :: String -- ^ name of the module + -> ByteString -- ^ NVVM IR in either textual or bitcode representation + -> [CompileOption] -- ^ compiler options + -> IO Result +compileModule !name !bs !opts = + compileModules [(name,bs)] opts + + +-- | Compile a collection of NVVM IR modules into PTX code +-- +{-# INLINEABLE compileModules #-} +compileModules + :: [(String, ByteString)] -- ^ (module name, module NVVM IR) pairs to compile + -> [CompileOption] -- ^ compiler options + -> IO Result +compileModules !bss !opts = + bracket create destroy $ \prg -> do + mapM_ (uncurry (addModule prg)) bss + (messages, result) <- compile prg opts + case result of + Nothing -> nvvmErrorIO (B.unpack messages) + Just ptx -> return $ Result ptx messages + + +-- | Create an empty 'Program' +-- +-- +-- +{-# INLINEABLE create #-} +create :: IO Program +create = resultIfOk =<< nvvmCreateProgram + where + peekProgram ptr = Program `fmap` peek ptr + {# + fun unsafe nvvmCreateProgram + { alloca- `Program' peekProgram* + } + -> `Status' cToEnum + #} + + +-- | Destroy a 'Program' +-- +-- +-- +{-# INLINEABLE destroy #-} +destroy :: Program -> IO () +destroy !prg = nothingIfOk =<< nvvmDestroyProgram prg + where + withProgram p = with (useProgram p) + {# + fun unsafe nvvmDestroyProgram + { withProgram* `Program' + } + -> `Status' cToEnum + #} + + +-- | Add a module level NVVM IR to a program +-- +-- +-- +{-# INLINEABLE addModule #-} +addModule + :: Program -- ^ NVVM program to add to + -> String -- ^ Name of the module (defaults to "" if empty) + -> ByteString -- ^ NVVM IR module in either bitcode or textual representation + -> IO () +addModule !prg !name !bs = + B.unsafeUseAsCStringLen bs $ \(ptr,len) -> + addModuleFromPtr prg name len (castPtr ptr) + + +-- | As with 'addModule', but read the specified number of bytes from the given +-- pointer. +-- +{-# INLINEABLE addModuleFromPtr #-} +addModuleFromPtr + :: Program -- ^ NVVM program to add to + -> String -- ^ Name of the module (defaults to "" if empty) + -> Int -- ^ Number of bytes in the module + -> Ptr Word8 -- ^ NVVM IR module in bitcode or textual representation + -> IO () +addModuleFromPtr !prg !name !size !buffer = + nothingIfOk =<< nvvmAddModuleToProgram prg buffer size name + where + {# + fun unsafe nvvmAddModuleToProgram + { useProgram `Program' + , castPtr `Ptr Word8' + , cIntConv `Int' + , withCString* `String' + } + -> `Status' cToEnum + #} + + +-- | Compile the NVVM program. Returns the log from the compiler/verifier and, +-- if successful, the compiled program. +-- +-- +-- +{-# INLINEABLE compile #-} +compile :: Program -> [CompileOption] -> IO (ByteString, Maybe ByteString) +compile !prg !opts = do + status <- withCompileOptions opts (nvvmCompileProgram prg) + messages <- retrieve (nvvmGetProgramLogSize prg) (nvvmGetProgramLog prg) + case status of + Success -> do ptx <- retrieve (nvvmGetCompiledResultSize prg) (nvvmGetCompiledResult prg) + return (messages, Just ptx) + _ -> return (messages, Nothing) + where + {# fun unsafe nvvmCompileProgram + { useProgram `Program' + , cIntConv `Int' + , id `Ptr CString' + } + -> `Status' cToEnum + #} + + {# fun unsafe nvvmGetCompiledResultSize + { useProgram `Program' + , alloca- `Int' peekIntConv* + } + -> `Status' cToEnum + #} + + {# fun unsafe nvvmGetCompiledResult + { useProgram `Program' + , withForeignPtr'* `ForeignPtr Word8' + } + -> `Status' cToEnum + #} + + +-- | Verify the NVVM program. Returns whether compilation will succeed, together +-- with any error or warning messages. +-- +{-# INLINEABLE verify #-} +verify :: Program -> [CompileOption] -> IO (Status, ByteString) +verify !prg !opts = do + status <- withCompileOptions opts (nvvmVerifyProgram prg) + messages <- retrieve (nvvmGetProgramLogSize prg) (nvvmGetProgramLog prg) + return (status, messages) + where + {# + fun unsafe nvvmVerifyProgram + { useProgram `Program' + , cIntConv `Int' + , id `Ptr CString' + } + -> `Status' cToEnum + #} + + +{# fun unsafe nvvmGetProgramLogSize + { useProgram `Program' + , alloca- `Int' peekIntConv* + } + -> `Status' cToEnum +#} + +{# fun unsafe nvvmGetProgramLog + { useProgram `Program' + , withForeignPtr'* `ForeignPtr Word8' + } + -> `Status' cToEnum +#} + + +-- Utilities +-- --------- + +{-# INLINEABLE withForeignPtr' #-} +withForeignPtr' :: ForeignPtr Word8 -> (Ptr CChar -> IO a) -> IO a +withForeignPtr' fp f = withForeignPtr fp (f . castPtr) + + +{-# INLINEABLE withCompileOptions #-} +withCompileOptions :: [CompileOption] -> (Int -> Ptr CString -> IO a) -> IO a +withCompileOptions opts next = + withMany withCString (map toStr opts) $ \cs -> withArrayLen cs next + where + toStr :: CompileOption -> String + toStr (OptimisationLevel n) = printf "-opt=%d" n + toStr (Target (Compute n m)) = printf "-arch=compute_%d%d" n m + toStr FlushToZero = "-ftz=1" + toStr NoFMA = "-fma=0" + toStr FastSqrt = "-prec-sqrt=0" + toStr FastDiv = "-prec-div=0" + toStr GenerateDebugInfo = "-g" + +{-# INLINEABLE retrieve #-} +retrieve :: IO (Status, Int) -> (ForeignPtr Word8 -> IO Status) -> IO ByteString +retrieve size payload = do + bytes <- resultIfOk =<< size + if bytes <= 1 -- size includes NULL terminator + then return B.empty + else do fp <- mallocForeignPtrBytes bytes + _ <- nothingIfOk =<< payload fp + return (B.fromForeignPtr fp 0 bytes) + diff --git a/Foreign/NVVM/Error.chs b/Foreign/NVVM/Error.chs new file mode 100644 --- /dev/null +++ b/Foreign/NVVM/Error.chs @@ -0,0 +1,114 @@ +{-# LANGUAGE DeriveDataTypeable #-} +-------------------------------------------------------------------------------- +-- | +-- Module : Foreign.NVVM.Error +-- Copyright : [2016] Trevor L. McDonell +-- License : BSD +-- +-- Error handling +-- +-------------------------------------------------------------------------------- + +module Foreign.NVVM.Error ( + + Status(..), + describe, + resultIfOk, nothingIfOk, + nvvmError, nvvmErrorIO, requireSDK, + +) where + +import Foreign.NVVM.Internal.C2HS +import Foreign.C.String + +import Control.Exception +import Data.Typeable +import Language.Haskell.TH +import Text.Printf + +#include "cbits/stubs.h" +{# context lib="nvvm" #} + + +-- Return codes +-- ------------ + +-- | NVVM API function return code +-- +{# enum nvvmResult as Status + { underscoreToCase + , NVVM_SUCCESS as Success + , NVVM_ERROR_IR_VERSION_MISMATCH as IRVersionMismatch + , NVVM_ERROR_INVALID_IR as InvalidIR + , NVVM_ERROR_COMPILATION as CompilationFailure + } + with prefix="NVVM_ERROR" deriving (Eq, Show) #} + + +-- | Get the descriptive message string for the given result code +-- +{# + fun pure unsafe nvvmGetErrorString as describe + { cFromEnum `Status' + } + -> `String' peekCString* +#} + + +-- Exceptions +-- ---------- + +data NVVMException + = ExitCode Status + | UserError String + deriving Typeable + +instance Exception NVVMException + +instance Show NVVMException where + showsPrec _ (ExitCode s) = showString ("NVVM Exception: " ++ describe s) + showsPrec _ (UserError s) = showString ("NVVM Exception: " ++ s) + + +-- | Throw an exception. Exceptions may be thrown from pure code, but can only +-- be caught in the 'IO' monad. +-- +{-# RULES "nvvmError/IO" nvvmError = nvvmErrorIO #-} +{-# NOINLINE [1] nvvmError #-} +nvvmError :: String -> a +nvvmError s = throw (UserError s) + +-- | Raise an NVVM exception in the 'IO' monad +-- +nvvmErrorIO :: String -> IO a +nvvmErrorIO s = throwIO (UserError s) + +-- | +-- A specially formatted error message +-- +requireSDK :: Name -> Double -> a +requireSDK n v = nvvmError $ printf "'%s' requires at least cuda-%3.1f\n" (show n) v + + +-- Helper functions +-- ---------------- + +-- | Return the result of a function on successful execution, otherwise throw an +-- exception. +-- +{-# INLINE resultIfOk #-} +resultIfOk :: (Status, a) -> IO a +resultIfOk (status, result) = + case status of + Success -> return $! result + _ -> throwIO (ExitCode status) + +-- | Throw an exception on an unsuccessful return code +-- +{-# INLINE nothingIfOk #-} +nothingIfOk :: Status -> IO () +nothingIfOk status = + case status of + Success -> return () + _ -> throwIO (ExitCode status) + diff --git a/Foreign/NVVM/Info.chs b/Foreign/NVVM/Info.chs new file mode 100644 --- /dev/null +++ b/Foreign/NVVM/Info.chs @@ -0,0 +1,80 @@ +{-# LANGUAGE CPP #-} +{-# LANGUAGE TemplateHaskell #-} +-------------------------------------------------------------------------------- +-- | +-- Module : Foreign.NVVM.Info +-- Copyright : [2016] Trevor L. McDonell +-- License : BSD +-- +-- General information query +-- +-------------------------------------------------------------------------------- + +module Foreign.NVVM.Info ( + + nvvmVersion, + nvvmIRVersion, + +) where + +import Foreign.NVVM.Error +import Foreign.NVVM.Internal.C2HS +import Foreign.Marshal + +import Data.Version +import System.IO.Unsafe + + +#include "cbits/stubs.h" +{# context lib="nvvm" #} + + +-- | Get the version of NVVM IR supported by this library. The first component +-- is the NVVM IR version, and the second the version of the debug metadata. +-- +-- Requires: CUDA-7.0 +-- +-- +-- +nvvmIRVersion :: (Version, Version) +#if CUDA_VERSION < 7000 +nvvmIRVersion = requireSDK 'nvvmIRVersion 7.0 +#else +nvvmIRVersion = unsafePerformIO $ do + (status, r1, r2, d1, d2) <- c_nvvmIRVersion + resultIfOk (status, (makeVersion [r1,r2], makeVersion [d1,d2])) + +{# + fun unsafe nvvmIRVersion as c_nvvmIRVersion + { alloca- `Int' peekIntConv* + , alloca- `Int' peekIntConv* + , alloca- `Int' peekIntConv* + , alloca- `Int' peekIntConv* + } + -> `Status' cToEnum +#} +#endif + + +-- | Get the version of the NVVM library +-- +-- +-- +nvvmVersion :: Version +nvvmVersion = unsafePerformIO $ do + (status, v1, v2) <- c_nvvmVersion + resultIfOk (status, makeVersion [v1,v2]) + +{# + fun unsafe nvvmVersion as c_nvvmVersion + { alloca- `Int' peekIntConv* + , alloca- `Int' peekIntConv* + } + -> `Status' cToEnum +#} + +#if __GLASGOW_HASKELL__ <= 708 +makeVersion :: [Int] -> Version +makeVersion vs = Version vs [] +#endif + diff --git a/Foreign/NVVM/Internal/C2HS.hs b/Foreign/NVVM/Internal/C2HS.hs new file mode 100644 --- /dev/null +++ b/Foreign/NVVM/Internal/C2HS.hs @@ -0,0 +1,230 @@ +-- C->Haskell Compiler: Marshalling library +-- +-- Copyright (c) [1999...2005] Manuel M T Chakravarty +-- +-- Redistribution and use in source and binary forms, with or without +-- modification, are permitted provided that the following conditions are met: +-- +-- 1. Redistributions of source code must retain the above copyright notice, +-- this list of conditions and the following disclaimer. +-- 2. Redistributions in binary form must reproduce the above copyright +-- notice, this list of conditions and the following disclaimer in the +-- documentation and/or other materials provided with the distribution. +-- 3. The name of the author may not be used to endorse or promote products +-- derived from this software without specific prior written permission. +-- +-- THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR +-- IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES +-- OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN +-- NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +-- SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED +-- TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +-- PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF +-- LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING +-- NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +-- SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +-- +--- Description --------------------------------------------------------------- +-- +-- Language: Haskell 98 +-- +-- This module provides the marshaling routines for Haskell files produced by +-- C->Haskell for binding to C library interfaces. It exports all of the +-- low-level FFI (language-independent plus the C-specific parts) together +-- with the C->HS-specific higher-level marshalling routines. +-- + +module Foreign.NVVM.Internal.C2HS ( + + -- * Composite marshalling functions + withCStringLenIntConv, peekCStringLenIntConv, withIntConv, withFloatConv, + peekIntConv, peekFloatConv, withBool, peekBool, withEnum, peekEnum, + peekArrayWith, + + -- * Conditional results using 'Maybe' + nothingIf, nothingIfNull, + + -- * Bit masks + combineBitMasks, containsBitMask, extractBitMasks, + + -- * Conversion between C and Haskell types + cIntConv, cFloatConv, cToBool, cFromBool, cToEnum, cFromEnum +) where + + +import Foreign +import Foreign.C +import Control.Monad ( liftM ) + + +-- Composite marshalling functions +-- ------------------------------- + +-- Strings with explicit length +-- +withCStringLenIntConv :: String -> (CStringLen -> IO a) -> IO a +withCStringLenIntConv s f = withCStringLen s $ \(p, n) -> f (p, cIntConv n) + +peekCStringLenIntConv :: CStringLen -> IO String +peekCStringLenIntConv (s, n) = peekCStringLen (s, cIntConv n) + +-- Marshalling of numerals +-- + +withIntConv :: (Storable b, Integral a, Integral b) => a -> (Ptr b -> IO c) -> IO c +withIntConv = with . cIntConv + +withFloatConv :: (Storable b, RealFloat a, RealFloat b) => a -> (Ptr b -> IO c) -> IO c +withFloatConv = with . cFloatConv + +peekIntConv :: (Storable a, Integral a, Integral b) => Ptr a -> IO b +peekIntConv = liftM cIntConv . peek + +peekFloatConv :: (Storable a, RealFloat a, RealFloat b) => Ptr a -> IO b +peekFloatConv = liftM cFloatConv . peek + +-- Passing Booleans by reference +-- + +withBool :: (Integral a, Storable a) => Bool -> (Ptr a -> IO b) -> IO b +withBool = with . fromBool + +peekBool :: (Integral a, Storable a) => Ptr a -> IO Bool +peekBool = liftM toBool . peek + + +-- Read and marshal array elements +-- + +peekArrayWith :: Storable a => (a -> b) -> Int -> Ptr a -> IO [b] +peekArrayWith f n p = map f `fmap` peekArray n p + + +-- Passing enums by reference +-- + +withEnum :: (Enum a, Integral b, Storable b) => a -> (Ptr b -> IO c) -> IO c +withEnum = with . cFromEnum + +peekEnum :: (Enum a, Integral b, Storable b) => Ptr b -> IO a +peekEnum = liftM cToEnum . peek + + +{- +-- Storing of 'Maybe' values +-- ------------------------- + +instance Storable a => Storable (Maybe a) where + sizeOf _ = sizeOf (undefined :: Ptr ()) + alignment _ = alignment (undefined :: Ptr ()) + + peek p = do + ptr <- peek (castPtr p) + if ptr == nullPtr + then return Nothing + else liftM Just $ peek ptr + + poke p v = do + ptr <- case v of + Nothing -> return nullPtr + Just v' -> new v' + poke (castPtr p) ptr +-} + + +-- Conditional results using 'Maybe' +-- --------------------------------- + +-- Wrap the result into a 'Maybe' type. +-- +-- * the predicate determines when the result is considered to be non-existing, +-- ie, it is represented by `Nothing' +-- +-- * the second argument allows to map a result wrapped into `Just' to some +-- other domain +-- +nothingIf :: (a -> Bool) -> (a -> b) -> a -> Maybe b +nothingIf p f x = if p x then Nothing else Just $ f x + +-- |Instance for special casing null pointers. +-- +nothingIfNull :: (Ptr a -> b) -> Ptr a -> Maybe b +nothingIfNull = nothingIf (== nullPtr) + + +-- Support for bit masks +-- --------------------- + +-- Given a list of enumeration values that represent bit masks, combine these +-- masks using bitwise disjunction. +-- +combineBitMasks :: (Enum a, Num b, Bits b) => [a] -> b +combineBitMasks = foldl (.|.) 0 . map (fromIntegral . fromEnum) + +-- Tests whether the given bit mask is contained in the given bit pattern +-- (i.e., all bits set in the mask are also set in the pattern). +-- +containsBitMask :: (Num a, Bits a, Enum b) => a -> b -> Bool +bits `containsBitMask` bm = let bm' = fromIntegral . fromEnum $ bm + in + bm' .&. bits == bm' + +-- |Given a bit pattern, yield all bit masks that it contains. +-- +-- * This does *not* attempt to compute a minimal set of bit masks that when +-- combined yield the bit pattern, instead all contained bit masks are +-- produced. +-- +extractBitMasks :: (Num a, Bits a, Enum b, Bounded b) => a -> [b] +extractBitMasks bits = + [bm | bm <- [minBound..maxBound], bits `containsBitMask` bm] + + +-- Conversion routines +-- ------------------- + +-- |Integral conversion +-- +{-# INLINE [1] cIntConv #-} +cIntConv :: (Integral a, Integral b) => a -> b +cIntConv = fromIntegral + +-- |Floating conversion +-- +{-# INLINE [1] cFloatConv #-} +cFloatConv :: (RealFloat a, RealFloat b) => a -> b +cFloatConv = realToFrac +-- As this conversion by default goes via `Rational', it can be very slow... +{-# RULES + "cFloatConv/Float->Float" forall (x::Float). cFloatConv x = x; + "cFloatConv/Double->Double" forall (x::Double). cFloatConv x = x; + "cFloatConv/Float->CFloat" forall (x::Float). cFloatConv x = CFloat x; + "cFloatConv/CFloat->Float" forall (x::Float). cFloatConv CFloat x = x; + "cFloatConv/Double->CDouble" forall (x::Double). cFloatConv x = CDouble x; + "cFloatConv/CDouble->Double" forall (x::Double). cFloatConv CDouble x = x + #-} + +-- |Obtain C value from Haskell 'Bool'. +-- +{-# INLINE [1] cFromBool #-} +cFromBool :: Num a => Bool -> a +cFromBool = fromBool + +-- |Obtain Haskell 'Bool' from C value. +-- +{-# INLINE [1] cToBool #-} +cToBool :: (Eq a, Num a) => a -> Bool +cToBool = toBool + +-- |Convert a C enumeration to Haskell. +-- +{-# INLINE [1] cToEnum #-} +cToEnum :: (Integral i, Enum e) => i -> e +cToEnum = toEnum . cIntConv + +-- |Convert a Haskell enumeration to C. +-- +{-# INLINE [1] cFromEnum #-} +cFromEnum :: (Enum e, Integral i) => e -> i +cFromEnum = cIntConv . fromEnum + diff --git a/LICENSE b/LICENSE new file mode 100644 --- /dev/null +++ b/LICENSE @@ -0,0 +1,30 @@ +Copyright (c) 2016, Trevor L. McDonell + +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + + * Redistributions of source code must retain the above copyright + notice, this list of conditions and the following disclaimer. + + * Redistributions in binary form must reproduce the above + copyright notice, this list of conditions and the following + disclaimer in the documentation and/or other materials provided + with the distribution. + + * Neither the name of Trevor L. McDonell nor the names of other + contributors may be used to endorse or promote products derived + from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS +"AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT +LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR +A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. diff --git a/README.md b/README.md new file mode 100644 --- /dev/null +++ b/README.md @@ -0,0 +1,30 @@ +Haskell FFI Bindings to NVVM +============================ + +[![Build status](https://travis-ci.org/tmcdonell/nvvm.svg?branch=master)](https://travis-ci.org/tmcdonell/nvvm) + +The NVVM library compiles [NVVM IR][nvvm-ir-spec] (a subset of LLVM IR) into PTX code which can +then be executed on NVIDIA GPUs. + +In contrast to the standard [NVPTX][nvptx-spec] target built in to the LLVM +toolchain, NVVM includes a set of proprietary optimisations which are otherwise +only available by compiling CUDA code with the `nvcc` compiler. On the other +hand, the version of LLVM that NVVM is internally based on typically lags the +public release by several generations (years), so these secret optimisations may +or may not be worthwhile to your application. + +The resulting PTX code can be loaded onto the GPU and executed using the [cuda +package][hs-cuda]. + +The NVVM library is a compiler component available a part of the CUDA toolkit: + + + +The configure step will look for your CUDA installation in the standard places, +and if the `nvcc` compiler is found in your `PATH`, relative to that. + + +[nvptx-spec]: http://llvm.org/docs/NVPTXUsage.html +[nvvm-ir-spec]: http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html +[hs-cuda]: https://github.com/tmcdonell/cuda + diff --git a/Setup.hs b/Setup.hs new file mode 100644 --- /dev/null +++ b/Setup.hs @@ -0,0 +1,523 @@ + +import Distribution.PackageDescription +import Distribution.PackageDescription.Parse +import Distribution.Simple +import Distribution.Simple.BuildPaths +import Distribution.Simple.Command +import Distribution.Simple.LocalBuildInfo +import Distribution.Simple.PreProcess hiding ( ppC2hs ) +import Distribution.Simple.Program +import Distribution.Simple.Program.Db +import Distribution.Simple.Setup +import Distribution.Simple.Utils +import Distribution.System +import Distribution.Verbosity + +import Control.Applicative +import Control.Exception +import Control.Monad +import System.Directory +import System.Environment +import System.FilePath +import System.IO.Error +import Text.Printf +import Prelude + + +-- Configuration +-- ------------- + +customBuildInfoFilePath :: FilePath +customBuildInfoFilePath = "nvvm" <.> "buildinfo" + +generatedBuldInfoFilePath :: FilePath +generatedBuldInfoFilePath = customBuildInfoFilePath <.> "generated" + + +-- Build setup +-- ----------- + +main :: IO () +main = defaultMainWithHooks customHooks + where + readHook get_verbosity args flags = do + noExtraFlags args + getHookedBuildInfo (fromFlag (get_verbosity flags)) + + -- Our readHook implementation uses our getHookedBuildInfo. We can't rely on + -- cabal's autoconfUserHooks since they don't handle user overwrites to + -- buildinfo like we do. + -- + customHooks = + simpleUserHooks + { preBuild = preBuildHook -- not using 'readHook' here because 'build' takes extra args + , preClean = readHook cleanVerbosity + , preCopy = readHook copyVerbosity + , preInst = readHook installVerbosity + , preHscolour = readHook hscolourVerbosity + , preHaddock = readHook haddockVerbosity + , preReg = readHook regVerbosity + , preUnreg = readHook regVerbosity + , postConf = postConfHook + , postBuild = postBuildHook + , hookedPreProcessors = ("chs", ppC2hs) : filter (\x -> fst x /= "chs") (hookedPreProcessors simpleUserHooks) + } + + -- The hook just loads the HookedBuildInfo generated by postConfHook, unless + -- there is user-provided info that overwrites it. + -- + preBuildHook :: Args -> BuildFlags -> IO HookedBuildInfo + preBuildHook _ flags = getHookedBuildInfo $ fromFlag $ buildVerbosity flags + + -- The hook scans system in search for CUDA Toolkit. If the toolkit is not + -- found, an error is raised. Otherwise the toolkit location is used to + -- create a `nvvm.buildinfo.generated` file with all the resulting flags. + postConfHook :: Args -> ConfigFlags -> PackageDescription -> LocalBuildInfo -> IO () + postConfHook args flags pkg_descr lbi = do + let + verbosity = fromFlag (configVerbosity flags) + currentPlatform = hostPlatform lbi + compilerId_ = compilerId (compiler lbi) + -- + noExtraFlags args + generateAndStoreBuildInfo verbosity currentPlatform compilerId_ generatedBuldInfoFilePath + validateLinker verbosity currentPlatform $ withPrograms lbi + -- + actualBuildInfoToUse <- getHookedBuildInfo verbosity + let pkg_descr' = updatePackageDescription actualBuildInfoToUse pkg_descr + postConf simpleUserHooks args flags pkg_descr' lbi + + -- This hook fixes the embedded LC_RPATHs in the generated .dylib on OSX. + postBuildHook :: Args -> BuildFlags -> PackageDescription -> LocalBuildInfo -> IO () + postBuildHook _ flags pkg_descr lbi = do + let + verbosity = fromFlag (buildVerbosity flags) + platform = hostPlatform lbi + cid = compilerId (compiler lbi) + uid = localUnitId lbi + sharedLib = buildDir lbi mkSharedLibName cid uid + Just extraLibDirs' = extraLibDirs . libBuildInfo <$> library pkg_descr + -- + updateLibraryRPATHs verbosity platform sharedLib extraLibDirs' + + +-- It seems that the GHC and/or Cabal developers don't quite understand how +-- dynamic linking works on OSX. Even though we have specified +-- '-optl-Wl,-rpath,...' as part of the configuration, this (sometimes?) gets +-- filtered out somewhere, and the resulting .dylib that is generated does not +-- have this path embedded as an LC_RPATH. The result is that the NVVM library +-- will not be found, resulting in a link-time error. +-- +-- On *nix (and versions of OSX previous to El Capitan 10.11), we could use +-- [DY]LD_LIBRARY_PATH to specify where to resolve @rpath locations, but that is +-- no longer an option on OSX due to System Integrity Protection. +-- +-- An alternate argument is that the CUDA installer should have updated the +-- install name (LC_ID_DYLIB) of its dynamic libraries to include the full +-- absolute path, rather than relying on @rpath in the first place, which is +-- what Apple's system libraries do for example. +-- +updateLibraryRPATHs :: Verbosity -> Platform -> FilePath -> [FilePath] -> IO () +updateLibraryRPATHs verbosity (Platform _ os) sharedLib extraLibDirs' = + when (os == OSX) $ do + exists <- doesFileExist sharedLib + unless exists $ die $ printf "Unexpected failure: library does not exist: %s" sharedLib + -- + mint <- findProgram verbosity "install_name_tool" + case mint of + Nothing -> notice verbosity $ "Could not locate 'install_name_tool' in order to update LC_RPATH entries. This is likely to cause problems later on." + Just install_name_tool -> + forM_ extraLibDirs' $ \libDir -> + runProgramInvocation verbosity $ simpleProgramInvocation install_name_tool ["-add_rpath", libDir, sharedLib] + + +-- Reads user-provided `nvvm.buildinfo` if present, otherwise loads `nvvm.buildinfo.generated` +-- Outputs message informing about the other possibility. +-- Calls die when neither of the files is available. +-- (generated one should be always present, as it is created in the post-conf step) +-- +getHookedBuildInfo :: Verbosity -> IO HookedBuildInfo +getHookedBuildInfo verbosity = do + customBuildInfoExists <- doesFileExist customBuildInfoFilePath + if customBuildInfoExists + then do + notice verbosity $ printf "The user-provided buildinfo from file '%s' will be used. To use default settings, delete this file." customBuildInfoFilePath + readHookedBuildInfo verbosity customBuildInfoFilePath + else do + generatedBuildInfoExists <- doesFileExist generatedBuldInfoFilePath + if generatedBuildInfoExists + then do + notice verbosity $ printf "Using build information from '%s'" generatedBuldInfoFilePath + notice verbosity $ printf "Provide a '%s' file to override this behaviour" customBuildInfoFilePath + readHookedBuildInfo verbosity generatedBuldInfoFilePath + else + die $ printf "Unexpected failure: neither the default '%s' nor custom '%s' exist" generatedBuldInfoFilePath customBuildInfoFilePath + + +-- Runs CUDA detection procedure and stores .buildinfo to a file. +-- +generateAndStoreBuildInfo :: Verbosity -> Platform -> CompilerId -> FilePath -> IO () +generateAndStoreBuildInfo verbosity platform (CompilerId _ghcFlavor ghcVersion) path = do + installPath <- findCUDAInstallPath verbosity platform + hbi <- libraryBuildInfo installPath platform ghcVersion + storeHookedBuildInfo verbosity path hbi + + +-- Try to locate CUDA installation by checking (in order): +-- +-- 1. CUDA_PATH environment variable +-- 2. Looking for `nvcc` in `PATH` +-- 3. Checking /usr/local/cuda +-- +-- In case of failure, calls die with the pretty long message from below. +-- +findCUDAInstallPath :: Verbosity -> Platform -> IO FilePath +findCUDAInstallPath verbosity platform = do + result <- findFirstValidLocation verbosity platform (candidateCUDAInstallPaths verbosity platform) + case result of + Just installPath -> do + notice verbosity $ printf "Found CUDA toolkit at: %s" installPath + return installPath + Nothing -> die cudaNotFoundMsg + + +-- Function iterates over action yielding possible locations, evaluating them +-- and returning the first valid one. Returns Nothing if no location matches. +-- +findFirstValidLocation :: Verbosity -> Platform -> [(IO FilePath, String)] -> IO (Maybe FilePath) +findFirstValidLocation verbosity platform = go + where + go :: [(IO FilePath, String)] -> IO (Maybe FilePath) + go [] = return Nothing + go (x:xs) = do + let (path,desc) = x + info verbosity $ printf "checking for %s" desc + found <- validateIOLocation verbosity platform path + if found + then Just `fmap` path + else go xs + + +-- Evaluates IO to obtain the path, handling any possible exceptions. +-- If path is evaluable and points to valid CUDA toolkit returns True. +-- +validateIOLocation :: Verbosity -> Platform -> IO FilePath -> IO Bool +validateIOLocation verbosity platform iopath = + let handler :: IOError -> IO Bool + handler err = do + info verbosity (show err) + return False + in + (iopath >>= validateLocation verbosity platform) `catch` handler + + +-- Checks whether given location looks like a valid CUDA toolkit directory by +-- testing whether the 'nvvm.h' file exists in the expected place. +-- +-- TODO: Ideally this should check also for 'libnvvm' and whether the library +-- exports relevant symbols. This should be achievable with some `nm` trickery. +-- +validateLocation :: Verbosity -> Platform -> FilePath -> IO Bool +validateLocation verbosity platform path = do + let nvvmHeader = nvvmIncludePath platform path "nvvm.h" + cudaHeader = cudaIncludePath platform path "cuda.h" + -- + exists <- (&&) <$> doesFileExist nvvmHeader <*> doesFileExist cudaHeader + info verbosity $ + if exists + then printf "Path accepted: %s" path + else printf "Path rejected: %s\nDoes not exist: %s or %s" path cudaHeader nvvmHeader + return exists + + +-- Returns pairs of (action yielding candidate path, String description of that location) +-- +candidateCUDAInstallPaths :: Verbosity -> Platform -> [(IO FilePath, String)] +candidateCUDAInstallPaths verbosity platform = + [ (getEnv "CUDA_PATH", "environment variable CUDA_PATH") + , (findInPath, "nvcc compiler executable in PATH") + , (return defaultPath, printf "default install location (%s)" defaultPath) + ] + where + findInPath :: IO FilePath + findInPath = do + nvccPath <- findProgramLocationOrError verbosity "nvcc" + -- The obtained path is likely TOOLKIT/bin/nvcc. We want to extract the + -- TOOLKIT part + return (takeDirectory $ takeDirectory nvccPath) + + defaultPath :: FilePath + defaultPath = defaultCUDAInstallPath platform + + +defaultCUDAInstallPath :: Platform -> FilePath +defaultCUDAInstallPath _ = "/usr/local/cuda" -- windows? + + +-- NOTE: this function throws an exception when there is no `nvcc` in PATH. +-- The exception contains a meaningful message. +-- +findProgramLocationOrError :: Verbosity -> String -> IO FilePath +findProgramLocationOrError verbosity execName = do + location <- findProgram verbosity execName + case location of + Just path -> return path + Nothing -> ioError $ mkIOError doesNotExistErrorType ("not found: " ++ execName) Nothing Nothing + +findProgram :: Verbosity -> FilePath -> IO (Maybe FilePath) +findProgram verbosity prog = do + result <- findProgramOnSearchPath verbosity defaultProgramSearchPath prog + case result of + Nothing -> return Nothing + Just (path,_) -> return (Just path) + + +cudaNotFoundMsg :: String +cudaNotFoundMsg = unlines + [ "********************************************************************************" + , "" + , "The configuration process failed to locate your CUDA installation. Ensure that you have installed both the developer driver and toolkit, available from:" + , "" + , "> http://developer.nvidia.com/cuda-downloads" + , "" + , "and make sure that `nvcc` is available in your PATH, or set the CUDA_PATH environment variable appropriately. Check the above output log and run the command directly to ensure it can be located." + , "" + , "If you have a non-standard installation, you can add additional search paths using --extra-include-dirs and --extra-lib-dirs. Note that 64-bit Linux flavours often require both `lib64` and `lib` library paths, in that order." + , "" + , "********************************************************************************" + ] + + +-- Generates build info with flags needed for CUDA Toolkit to be properly +-- visible to underlying build tools. +-- +libraryBuildInfo :: FilePath -> Platform -> Version -> IO HookedBuildInfo +libraryBuildInfo installPath platform@(Platform arch os) ghcVersion = do + let + libraryPaths = [nvvmLibraryPath platform installPath] + includePaths = [nvvmIncludePath platform installPath, cudaIncludePath platform installPath] + + -- options for GHC + extraLibDirs' = libraryPaths + ccOptions' = map ("-I"++) includePaths + ldOptions' = map ("-L"++) extraLibDirs' + ghcOptions = map ("-optc"++) ccOptions' + ++ map ("-optl"++) ldOptions' + ++ if os /= Windows + then map ("-optl-Wl,-rpath,"++) extraLibDirs' + else [] + extraLibs' = nvvmLibrary platform + + -- options for C2HS + archFlag = case arch of + I386 -> "-m32" + X86_64 -> "-m64" + _ -> "" + emptyCase = ["-DUSE_EMPTY_CASE" | versionBranch ghcVersion >= [7,8]] + blocksExtension = ["-U__BLOCKS__" | os == OSX ] + c2hsOptions = unwords $ map ("--cppopts="++) ("-E" : archFlag : emptyCase ++ blocksExtension) + c2hsExtraOptions = ("x-extra-c2hs-options", c2hsOptions) + + addSystemSpecificOptions :: BuildInfo -> IO BuildInfo + addSystemSpecificOptions bi = + case os of + -- In the CUDA package this is used to populate the extraGHCiLibs + -- field with the mangled .dll names. I'm not sure what those are for + -- this library, so left out for the time being. + Windows -> return bi + OSX -> return bi + _ -> return bi + + buildInfo' <- addSystemSpecificOptions $ emptyBuildInfo + { ccOptions = ccOptions' + , ldOptions = ldOptions' + , extraLibs = extraLibs' + , extraLibDirs = extraLibDirs' + , options = [(GHC, ghcOptions) | os /= Windows] + , customFieldsBI = [c2hsExtraOptions] + } + + return (Just buildInfo', []) + + +-- Location of the NVVM library relative to the base CUDA installation +-- +nvvmPath :: Platform -> FilePath -> FilePath +nvvmPath _ base = base "nvvm" + +nvvmIncludePath :: Platform -> FilePath -> FilePath +nvvmIncludePath platform base = nvvmPath platform base "include" + +cudaIncludePath :: Platform -> FilePath -> FilePath +cudaIncludePath _ base = base "include" + +nvvmLibraryPath :: Platform -> FilePath -> FilePath +nvvmLibraryPath platform@(Platform arch os) base = nvvmPath platform base libpath + where + libpath = + case (os, arch) of + (Windows, I386) -> "Win32" + (Windows, X86_64) -> "x64" + (OSX, _) -> "lib" -- MacOS does not distinguish 32- vs. 64-bit paths + (_, X86_64) -> "lib64" -- treat all others similarly + _ -> "lib" + +nvvmLibrary :: Platform -> [String] +nvvmLibrary _ = ["nvvm"] + +storeHookedBuildInfo :: Verbosity -> FilePath -> HookedBuildInfo -> IO () +storeHookedBuildInfo verbosity path hbi = do + notice verbosity $ "Storing parameters to " ++ path + writeHookedBuildInfo path hbi + + +-- On Windows platform the binutils linker targeting x64 is bugged and cannot +-- properly link with import libraries generated by MS compiler (like the CUDA ones). +-- The programs would correctly compile and crash as soon as the first FFI call is made. +-- +-- Therefore we fail configure process if the linker is too old and provide user +-- with guidelines on how to fix the problem. +-- +validateLinker :: Verbosity -> Platform -> ProgramDb -> IO () +validateLinker verbosity (Platform X86_64 Windows) db = do + let say msg = printf "%s. If generated executables crash when making calls to NVVM please see: %s" msg windowsHelpPage + -- + maybeLdPath <- getRealLdPath verbosity db + case maybeLdPath of + Nothing -> warn verbosity $ say "Cannot find ld.exe to check if it is new enough" + Just ldPath -> do + debug verbosity $ "Checking if ld.exe at " ++ ldPath ++ " is new enough" + maybeVersion <- getLdVersion verbosity ldPath + case maybeVersion of + Nothing -> warn verbosity $ say "Unknown ld.exe version" + Just ldVersion -> do + debug verbosity $ "Found ld.exe version: " ++ show ldVersion + when (ldVersion < [2,25,1]) $ die (windowsLinkerBugMsg ldPath) +validateLinker _ _ _ = return () -- The linker bug is present only on Win64 platform + + +-- On Windows GHC package comes with two copies of ld.exe. +-- +-- ProgramDb knows about the first one: ghcpath\mingw\bin\ld.exe +-- This function returns the other one: ghcpath\mingw\x86_64-w64-mingw32\bin\ld.exe +-- +-- The second one is the one that does actual linking and code generation. +-- See: https://github.com/tmcdonell/cuda/issues/31#issuecomment-149181376 +-- +-- The function is meant to be used only on 64-bit GHC distributions. +-- +getRealLdPath :: Verbosity -> ProgramDb -> IO (Maybe FilePath) +getRealLdPath verbosity programDb = + -- This should ideally work `programFindVersion ldProgram` but for some reason + -- it does not. The issue should be investigated at some time. + case lookupProgram ghcProgram programDb of + Nothing -> return Nothing + Just configuredGhc -> do + let ghcPath = locationPath $ programLocation configuredGhc + presumedLdPath = (takeDirectory . takeDirectory) ghcPath "mingw" "x86_64-w64-mingw32" "bin" "ld.exe" + info verbosity $ "Presuming ld location" ++ presumedLdPath + presumedLdExists <- doesFileExist presumedLdPath + return $ if presumedLdExists + then Just presumedLdPath + else Nothing + +-- Tries to obtain the version `ld`. Throws an exception on failure. +-- +getLdVersion :: Verbosity -> FilePath -> IO (Maybe [Int]) +getLdVersion verbosity ldPath = do + -- Examples of version string format: + -- * GNU ld (GNU Binutils) 2.25.1 + -- * GNU ld (GNU Binutils) 2.20.51.20100613 + -- + ldVersionString <- getProgramInvocationOutput normal (simpleProgramInvocation ldPath ["-v"]) + + let versionText = last $ words ldVersionString -- takes e.g. "2.25.1" + versionParts = splitOn (== '.') versionText + versionParsed = Just $ map read versionParts + + -- last and read above may throw and message would be not understandable for user, + -- so we'll intercept exception and rethrow it with more useful message. + let handleError :: SomeException -> IO (Maybe [Int]) + handleError e = do + warn verbosity $ printf "cannot parse ld version string: '%s'. Parsing exception: %s" ldVersionString (show e) + return Nothing + + evaluate versionParsed `catch` handleError + +splitOn :: (Char -> Bool) -> String -> [String] +splitOn p s = + case dropWhile p s of + [] -> [] + ss -> let (w,s') = break p ss in w : splitOn p s' + + +windowsHelpPage :: String +windowsHelpPage = "https://github.com/tmcdonell/nvvm/blob/master/WINDOWS.md" + +windowsLinkerBugMsg :: FilePath -> String +windowsLinkerBugMsg ldPath = printf (unlines msg) windowsHelpPage ldPath + where + msg = + [ "********************************************************************************" + , "" + , "The installed version of `ld.exe` has version < 2.25.1. This version has known bug on Windows x64 architecture, making it unable to correctly link programs using CUDA. The fix is available and MSys2 released fixed version of `ld.exe` as part of their binutils package (version 2.25.1)." + , "" + , "To fix this issue, replace the `ld.exe` in your GHC installation with the correct binary. See the following page for details:" + , "" + , " %s" + , "" + , "The full path to the outdated `ld.exe` detected in your installation:" + , "" + , "> %s" + , "" + , "Please download a recent version of binutils `ld.exe`, from, e.g.:" + , "" + , " http://repo.msys2.org/mingw/x86_64/mingw-w64-x86_64-binutils-2.25.1-1-any.pkg.tar.xz" + , "" + , "********************************************************************************" + ] + + +-- Replicate the default C2HS preprocessor hook here, and inject a value for +-- extra-c2hs-options, if it was present in the buildinfo file +-- +-- Everything below copied from Distribution.Simple.PreProcess +-- +ppC2hs :: BuildInfo -> LocalBuildInfo -> PreProcessor +ppC2hs bi lbi = + PreProcessor + { platformIndependent = False + , runPreProcessor = \(inBaseDir, inRelativeFile) (outBaseDir, outRelativeFile) verbosity -> + rawSystemProgramConf verbosity c2hsProgram (withPrograms lbi) . filter (not . null) $ + maybe [] words (lookup "x-extra-c2hs-options" (customFieldsBI bi)) + ++ ["--include=" ++ outBaseDir] + ++ ["--cppopts=" ++ opt | opt <- getCppOptions bi lbi] + ++ ["--output-dir=" ++ outBaseDir, + "--output=" ++ outRelativeFile, + inBaseDir inRelativeFile] + } + +getCppOptions :: BuildInfo -> LocalBuildInfo -> [String] +getCppOptions bi lbi + = hcDefines (compiler lbi) + ++ ["-I" ++ dir | dir <- includeDirs bi] + ++ [opt | opt@('-':c:_) <- ccOptions bi, c `elem` "DIU"] + +hcDefines :: Compiler -> [String] +hcDefines comp = + case compilerFlavor comp of + GHC -> ["-D__GLASGOW_HASKELL__=" ++ versionInt version] + JHC -> ["-D__JHC__=" ++ versionInt version] + NHC -> ["-D__NHC__=" ++ versionInt version] + Hugs -> ["-D__HUGS__"] + _ -> [] + where version = compilerVersion comp + +-- TODO: move this into the compiler abstraction +-- FIXME: this forces GHC's crazy 4.8.2 -> 408 convention on all the other +-- compilers. Check if that's really what they want. +versionInt :: Version -> String +versionInt (Version { versionBranch = [] }) = "1" +versionInt (Version { versionBranch = [n] }) = show n +versionInt (Version { versionBranch = n1:n2:_ }) = printf "%d%02d" n1 n2 + diff --git a/cbits/stubs.h b/cbits/stubs.h new file mode 100644 --- /dev/null +++ b/cbits/stubs.h @@ -0,0 +1,18 @@ +/******************************************************************************* + * + * Module : stubs + * Copyright : [2016] Trevor L. McDonell + * License : BSD + * + * Extra bits for NVVM bindings + * + ******************************************************************************/ + +#ifndef C_STUBS_H +#define C_STUBS_H + +#include +#include // required for CUDA_VERSION, a proxy for NVVM_VERSION + +#endif /* C_STUBS_H */ + diff --git a/nvvm.cabal b/nvvm.cabal new file mode 100644 --- /dev/null +++ b/nvvm.cabal @@ -0,0 +1,81 @@ +name: nvvm +version: 0.7.5.0 +synopsis: FFI bindings to NVVM +description: + The NVVM library compiles NVVM IR (a subset of LLVM IR) into PTX code which + can then be executed on NVIDIA GPUs. In contrast to the standard NVPTX target + built in to the LLVM toolchain, NVVM includes a set of proprietary + optimisations which are otherwise only available by compiling CUDA code with + the `nvcc` compiler. + . + The resulting PTX code can be loaded onto the GPU and executed using the + 'cuda' package: + . + + . + The NVVM library is a compiler component available a part of the CUDA toolkit: + . + + . + The configure step will look for your CUDA installation in the standard + places, and if the `nvcc` compiler is found in your `PATH`, relative to that. + . + This package tested with version 7.5 of the CUDA toolkit. + . + +license: BSD3 +license-file: LICENSE +homepage: https://github.com/tmcdonell/nvvm +author: Trevor L. McDonell +maintainer: Trevor L. McDonell +copyright: [2016] Trevor L. McDonell +category: Foreign +build-type: Custom +cabal-version: >= 1.24 + +extra-source-files: + CHANGELOG.md + README.md + cbits/stubs.h + +extra-tmp-files: + nvvm.buildinfo.generated + +custom-setup + setup-depends: + base >= 4.6 + , Cabal >= 1.24 + , directory >= 1.0 + , filepath >= 1.0 + +library + default-language: Haskell2010 + include-dirs: . + ghc-options: -Wall -O2 + + exposed-modules: + Foreign.NVVM + Foreign.NVVM.Compile + Foreign.NVVM.Error + Foreign.NVVM.Info + + other-modules: + Foreign.NVVM.Internal.C2HS + + build-depends: + base >= 4.6 && < 5 + , bytestring + , cuda >= 0.7 + , template-haskell + + build-tools: + c2hs >= 0.21 + + +source-repository this + type: git + location: https://github.com/tmcdonell/nvvm + tag: 0.7.5.0 + +-- vim: nospell +