Skip to content

HTTPS clone URL

Subversion checkout URL

You can clone with HTTPS or Subversion.

Download ZIP
Browse files

test: added idle test program

  • Loading branch information...
commit d6932580214cd3e38a4b8f785c7cad46059e9583 1 parent ef448bd
Shinpei Kato authored
2  test/cuda/common/fmadd.c
View
@@ -79,7 +79,7 @@ int cuda_test_fmadd(unsigned int n, char *path)
return -1;
}
- res = cuDeviceGet(&dev, 0);
+ res = cuDeviceGet(&dev, 15);
if (res != CUDA_SUCCESS) {
printf("cuDeviceGet failed: res = %lu\n", (unsigned long)res);
return -1;
2  test/cuda/common/fmmul.c
View
@@ -79,7 +79,7 @@ int cuda_test_fmmul(unsigned int n, char *path)
return -1;
}
- res = cuDeviceGet(&dev, 0);
+ res = cuDeviceGet(&dev, 3);
if (res != CUDA_SUCCESS) {
printf("cuDeviceGet failed: res = %lu\n", (unsigned long)res);
return -1;
149 test/cuda/common/idle.c
View
@@ -0,0 +1,149 @@
+#include <cuda.h>
+#ifdef __KERNEL__ /* just for measurement */
+#include <linux/vmalloc.h>
+#include <linux/time.h>
+#define printf printk
+#define malloc vmalloc
+#define free vfree
+#define gettimeofday(x, y) do_gettimeofday(x)
+#else /* just for measurement */
+#include <sys/time.h>
+#include <stdio.h>
+#include <stdlib.h>
+#endif
+
+/* tvsub: ret = x - y. */
+static inline void tvsub(struct timeval *x,
+ struct timeval *y,
+ struct timeval *ret)
+{
+ ret->tv_sec = x->tv_sec - y->tv_sec;
+ ret->tv_usec = x->tv_usec - y->tv_usec;
+ if (ret->tv_usec < 0) {
+ ret->tv_sec--;
+ ret->tv_usec += 1000000;
+ }
+}
+
+int cuda_test_idle(unsigned int n, char *path)
+{
+ int i, j, idx;
+ CUresult res;
+ CUdevice dev;
+ CUcontext ctx;
+ CUfunction function;
+ CUmodule module;
+ CUdeviceptr d_data;
+ int block_x, block_y, grid_x, grid_y;
+ char fname[256];
+ struct timeval tv;
+ struct timeval tv_exec_start, tv_exec_end;
+ float exec;
+
+ block_x = 1;
+ block_y = 1;
+ grid_x = 1;
+ grid_y = 1;
+
+ res = cuInit(0);
+ if (res != CUDA_SUCCESS) {
+ printf("cuInit failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ res = cuDeviceGet(&dev, 0);
+ if (res != CUDA_SUCCESS) {
+ printf("cuDeviceGet failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ res = cuCtxCreate(&ctx, 0, dev);
+ if (res != CUDA_SUCCESS) {
+ printf("cuCtxCreate failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ sprintf(fname, "%s/idle_gpu.cubin", path);
+ res = cuModuleLoad(&module, fname);
+ if (res != CUDA_SUCCESS) {
+ printf("cuModuleLoad() failed\n");
+ return -1;
+ }
+ res = cuModuleGetFunction(&function, module, "idle");
+ if (res != CUDA_SUCCESS) {
+ printf("cuModuleGetFunction() failed\n");
+ return -1;
+ }
+ res = cuFuncSetSharedSize(function, 0);
+ if (res != CUDA_SUCCESS) {
+ printf("cuFuncSetSharedSize() failed\n");
+ return -1;
+ }
+ res = cuFuncSetBlockShape(function, block_x, block_y, 1);
+ if (res != CUDA_SUCCESS) {
+ printf("cuFuncSetBlockShape() failed\n");
+ return -1;
+ }
+
+ res = cuMemAlloc(&d_data, sizeof(unsigned int));
+
+ /* set kernel parameters */
+ res = cuParamSeti(function, 0, d_data);
+ if (res != CUDA_SUCCESS) {
+ printf("cuParamSeti failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ res = cuParamSeti(function, 4, d_data>>32);
+ if (res != CUDA_SUCCESS) {
+ printf("cuParamSeti failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ res = cuParamSeti(function, 8, n);
+ if (res != CUDA_SUCCESS) {
+ printf("cuParamSeti failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ res = cuParamSetSize(function, 12);
+ if (res != CUDA_SUCCESS) {
+ printf("cuParamSetSize failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ printf("n = %u\n", n);
+ gettimeofday(&tv_exec_start, NULL);
+ res = cuLaunchGrid(function, grid_x, grid_y);
+ if (res != CUDA_SUCCESS) {
+ printf("cuLaunchGrid failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+ cuCtxSynchronize();
+ gettimeofday(&tv_exec_end, NULL);
+
+ cuMemcpyDtoH(&n, d_data, sizeof(n));
+
+ printf("n = %u\n", n);
+
+ cuMemFree(d_data);
+
+ res = cuModuleUnload(module);
+ if (res != CUDA_SUCCESS) {
+ printf("cuModuleUnload failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ res = cuCtxDestroy(ctx);
+ if (res != CUDA_SUCCESS) {
+ printf("cuCtxDestroy failed: res = %lu\n", (unsigned long)res);
+ return -1;
+ }
+
+ tvsub(&tv_exec_end, &tv_exec_start, &tv);
+ exec = tv.tv_sec * 1000.0 + (float) tv.tv_usec / 1000.0;
+
+ printf("Exec: %f\n", exec);
+
+ return 0;
+}
24 test/cuda/common/idle_gpu.cu
View
@@ -0,0 +1,24 @@
+#include <stdint.h>
+#include <cuda.h>
+
+extern "C"
+__global__
+void idle(unsigned int *p, unsigned int n)
+{
+ int x = blockIdx.x * blockDim.x + threadIdx.x;
+ int y = blockIdx.y * blockDim.y + threadIdx.y;
+ unsigned int i = 0, j = 0, k = 0;
+ __shared__ int s;
+
+ s = *p;
+ if (x == 0 && y == 0) {
+ for (i = 0; i < n; i++) {
+ if (x + y > n) {
+ s = s + x;
+ if (s > x + y)
+ s = x;
+ }
+ }
+ }
+ *p = s;
+}
13 test/cuda/user/idle/Makefile
View
@@ -0,0 +1,13 @@
+# Makefile
+TARGET = user_test
+CC = gcc
+NVCC = nvcc -arch sm_20 -cubin
+LIBS = -lcuda -lgdev
+CFLAGS = -L /usr/local/gdev/lib64 -I /usr/local/gdev/include
+
+all:
+ $(NVCC) -o idle_gpu.cubin idle_gpu.cu
+ gcc -o $(TARGET) $(CFLAGS) $(LIBS) main.c idle.c
+
+clean:
+ rm -f $(TARGET) *.cubin ./*~
13 test/cuda/user/idle/Makefile.nvidia
View
@@ -0,0 +1,13 @@
+# Makefile
+TARGET = user_test
+CC = gcc
+NVCC = nvcc -arch sm_20 -cubin
+LIBS = -lcuda
+CFLAGS = -I /usr/local/cuda/include
+
+all:
+ $(NVCC) -o idle_gpu.cubin idle_gpu.cu
+ gcc -o $(TARGET) $(CFLAGS) $(LIBS) main.c idle.c
+
+clean:
+ rm -f $(TARGET) *.cubin ./*~
1  test/cuda/user/idle/idle.c
View
1  test/cuda/user/idle/idle_gpu.cu
View
18 test/cuda/user/idle/main.c
View
@@ -0,0 +1,18 @@
+#include <stdio.h>
+
+int cuda_test_idle(unsigned int n, char *path);
+
+int main(int argc, char *argv[])
+{
+ unsigned int n = 10000000;
+
+ if (argc > 1)
+ n = atoi(argv[1]);
+
+ if (cuda_test_idle(n, ".") < 0)
+ printf("Test failed\n");
+ else
+ printf("Test passed\n");
+
+ return 0;
+}
Please sign in to comment.
Something went wrong with that request. Please try again.