nvvm (empty) → 0.7.5.0
raw patch · 11 files changed
+1584/−0 lines, 11 filesdep +basedep +bytestringdep +cudabuild-type:Customsetup-changed
Dependencies added: base, bytestring, cuda, template-haskell
Files
- CHANGELOG.md +5/−0
- Foreign/NVVM.hs +190/−0
- Foreign/NVVM/Compile.chs +283/−0
- Foreign/NVVM/Error.chs +114/−0
- Foreign/NVVM/Info.chs +80/−0
- Foreign/NVVM/Internal/C2HS.hs +230/−0
- LICENSE +30/−0
- README.md +30/−0
- Setup.hs +523/−0
- cbits/stubs.h +18/−0
- nvvm.cabal +81/−0
+ CHANGELOG.md view
@@ -0,0 +1,5 @@+# Revision history for nvvm++## 0.7.5.0 -- YYYY-mm-dd++* First version. Released on an unsuspecting world.
+ Foreign/NVVM.hs view
@@ -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:+--+-- <http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html>+--+-- 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:+--+-- <https://hackage.haskell.org/package/cuda>+--+--+-- [/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:+--+-- <http://docs.nvidia.com/cuda/libdevice-users-guide/index.html>+--+--------------------------------------------------------------------------------++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+
+ Foreign/NVVM/Compile.chs view
@@ -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'+--+-- <http://docs.nvidia.com/cuda/libnvvm-api/group__compilation.html#group__compilation_1g46a0ab04a063cba28bfbb41a1939e3f4>+--+{-# 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'+--+-- <http://docs.nvidia.com/cuda/libnvvm-api/group__compilation.html#group__compilation_1gfba94cab1224c0152841b80690d366aa>+--+{-# 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+--+-- <http://docs.nvidia.com/cuda/libnvvm-api/group__compilation.html#group__compilation_1g0c22d2b9be033c165bc37b16f3ed75c6>+--+{-# INLINEABLE addModule #-}+addModule+ :: Program -- ^ NVVM program to add to+ -> String -- ^ Name of the module (defaults to "<unnamed>" 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 "<unnamed>" 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.+--+-- <http://docs.nvidia.com/cuda/libnvvm-api/group__compilation.html#group__compilation_1g76ac1e23f5d0e2240e78be0e63450346>+--+{-# 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)+
+ Foreign/NVVM/Error.chs view
@@ -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)+
+ Foreign/NVVM/Info.chs view
@@ -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+--+-- <http://docs.nvidia.com/cuda/libnvvm-api/group__query.html#group__query_1g0894677934db095b3c40d4f8e2578cc5>+--+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+--+-- <http://docs.nvidia.com/cuda/libnvvm-api/group__query.html#group__query_1gcdd062f26078d20ded68f1017e999246>+--+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+
+ Foreign/NVVM/Internal/C2HS.hs view
@@ -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+
+ LICENSE view
@@ -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.
+ README.md view
@@ -0,0 +1,30 @@+Haskell FFI Bindings to NVVM+============================++[](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:++ <https://developer.nvidia.com/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+
+ Setup.hs view
@@ -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+
+ cbits/stubs.h view
@@ -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 <nvvm.h>+#include <cuda.h> // required for CUDA_VERSION, a proxy for NVVM_VERSION++#endif /* C_STUBS_H */+
+ nvvm.cabal view
@@ -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:+ .+ <https://hackage.haskell.org/package/cuda>+ .+ The NVVM library is a compiler component available a part of the CUDA toolkit:+ .+ <https://developer.nvidia.com/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 <tmcdonell@cse.unsw.edu.au>+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+