From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001
From: Tor Aamodt <aamodt@ece.ubc.ca>
Date: Thu, 15 Jul 2010 18:09:46 -0800
Subject: creating branch for adding support for CUDA 3.x and Fermi

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829]
---
 benchmarks/CUDA/WP/util4.cu | 46 +++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 46 insertions(+)
 create mode 100644 benchmarks/CUDA/WP/util4.cu

(limited to 'benchmarks/CUDA/WP/util4.cu')

diff --git a/benchmarks/CUDA/WP/util4.cu b/benchmarks/CUDA/WP/util4.cu
new file mode 100644
index 0000000..4b629d4
--- /dev/null
+++ b/benchmarks/CUDA/WP/util4.cu
@@ -0,0 +1,46 @@
+#ifndef PREPASS
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include "cublas.h"
+#endif
+
+__device__ float4 max4 ( const float4 a , const float4 b )
+{
+    float4 c  ;
+    c.x = (a.x>b.x)?a.x:b.x;
+    c.y = (a.y>b.y)?a.y:b.y;
+    c.z = (a.z>b.z)?a.z:b.z;
+    c.w = (a.w>b.w)?a.w:b.w;
+    return(c) ;
+}
+__device__ float4 min4 ( const float4 a , const float4 b )
+{
+    float4 c  ;
+    c.x = (a.x<b.x)?a.x:b.x;
+    c.y = (a.y<b.y)?a.y:b.y;
+    c.z = (a.z<b.z)?a.z:b.z;
+    c.w = (a.w<b.w)?a.w:b.w;
+    return(c) ;
+}
+
+__device__ float4 log4 ( const float4 a )
+{
+    float4 c  ;
+    c.x = log(a.x) ; c.y = log(a.y) ; c.z = log(a.z) ; c.w = log(a.w) ;
+    return(c) ;
+}
+
+__device__ float4 exp4 ( const float4 a )
+{
+    float4 c  ;
+    c.x = exp(a.x) ; c.y = exp(a.y) ; c.z = exp(a.z) ; c.w = exp(a.w) ;
+    return(c) ;
+}
+
+__device__ float4 sqrt4 ( const float4 a )
+{
+    float4 c  ;
+    c.x = sqrt(a.x) ; c.y = sqrt(a.y) ; c.z = sqrt(a.z) ; c.w = sqrt(a.w) ;
+    return(c) ;
+}
-- 
cgit v1.3