changeset 283:2b41bd298fe8

add openCL test files
author mir3636
date Sun, 05 Feb 2017 18:30:30 +0900
parents a3448b0f0a56
children e6bc0a4c2c36
files src/parallel_execution/origin_cs.c src/parallel_execution/origin_cs.h src/test/Cudasample_cpu.cu src/test/Cudasample_gpu.cu src/test/OpenCL_gpu.c
diffstat 5 files changed, 189 insertions(+), 31 deletions(-) [+]
line wrap: on
line diff
--- a/src/parallel_execution/origin_cs.c	Sun Feb 05 04:08:30 2017 +0900
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,27 +0,0 @@
-#include <stdlib.h>
-#include <stdio.h>
-#include "context.h"
-
-__code meta(struct Context* context, enum Code next) {
-    // printf("meta %d\n",next);
-    goto (context->code[next])(context);
-}
-
-__code start_code(struct Context* context) {
-    goto meta(context, context->next);
-}
-
-__code start_code_stub(struct Context* context) {
-    goto start_code(context);
-}
-
-__code exit_code(struct Context* context) {
-    free(context->code);
-    free(context->data);
-    free(context->heapStart);
-    goto exit(0);
-}
-
-__code exit_code_stub(struct Context* context) {
-    goto exit_code(context);
-}
--- a/src/parallel_execution/origin_cs.h	Sun Feb 05 04:08:30 2017 +0900
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,4 +0,0 @@
-extern __code start_code(struct Context* context);
-extern __code exit_code(struct Context* context);
-extern __code meta(struct Context* context, enum Code next);
-extern void initContext(struct Context* context);
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/test/Cudasample_cpu.cu	Sun Feb 05 18:30:30 2017 +0900
@@ -0,0 +1,26 @@
+#include <stdio.h>
+
+int main(void)
+{       
+        int b;
+      
+        for (b = 99; b >= 0; b--) {
+                switch (b) {
+                case 0:
+                        printf("No more bottles of beer on the wall, no more bottles of beer.\n");
+                        printf("Go to the store and buy some more, 99 bottles of beer on the wall.\n");
+                        break;
+                case 1:
+                        printf("1 bottle of beer on the wall, 1 bottle of beer.\n");
+                        printf("Take one down and pass it around, no more bottles of beer on the wall\n");
+                        break;
+                default:
+                        printf("%d bottles of beer on the wall, %d bottles of beer.\n", b, b);
+                        printf("Take one down and pass it around, %d %s of beer on the wall.\n"
+                                ,b - 1
+                                ,((b - 1) > 1)? "bottles" : "bottle");
+                        break;
+                }
+        }               
+        return 0;
+}      
\ No newline at end of file
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/test/Cudasample_gpu.cu	Sun Feb 05 18:30:30 2017 +0900
@@ -0,0 +1,72 @@
+#include <stdio.h>
+
+#define SIZE_TEXT (sizeof(text)-1)
+#define SIZE_END (sizeof(end)-1)
+
+__device__ char text[] =
+"__ bottles of beer on the wall, __ bottles of beer!\n"
+"Take one down, and pass it around, ## bottles of beer on the wall!\n\n";
+       
+__device__ char end[] =
+"01 bottle of beer on the wall, 01 bottle of beer.\n"
+"Take one down and pass it around, no more bottles of beer on the wall.\n"
+"\n"
+"No more bottles of beer on the wall, no more bottles of beer.\n"
+"Go to the store and buy some more, 99 bottles of beer on the wall.";
+
+
+__global__
+void bottle99(char *addr){
+    int x = threadIdx.x;
+       addr += x * SIZE_TEXT;
+    int bottle = 99 - x;
+       if (bottle == 1) {
+        for (int i=0; i<SIZE_END; i++) {
+               addr[i] = end[i];
+           }
+         addr[SIZE_END] = '\0';
+           } else {
+               char c1 = (bottle/10) + '0';
+               char c2 = (bottle%10) + '0';
+
+               char d1 = ((bottle-1)/10) + '0';
+               char d2 = ((bottle-1)%10) + '0';
+      
+           for (int i=0; i<SIZE_TEXT; i++) {
+               int c = text[i];
+               if (c == '_') {
+                  addr[i] = c1;
+                  addr[i+1] = c2;
+                  i++;
+               } else if (c == '#') {
+
+            addr[i] = d1;
+               addr[i+1] = d2;
+                    i++;
+               } else {
+
+                 addr[i] = text[i];
+            }
+        }
+    }
+}
+      
+int main()
+{
+    char *buffer;
+    char *d_buffer;
+      
+    int size = SIZE_TEXT * 98 + SIZE_END + 1;
+
+    buffer = new char[size];
+    cudaMalloc((void**)&d_buffer, size);
+
+    bottle99<<<1, 99>>>(d_buffer);
+
+    cudaMemcpy(buffer, d_buffer, size, cudaMemcpyDeviceToHost);
+    cudaFree(d_buffer);
+
+    puts(buffer);
+    free(buffer);               
+       }
+  
\ No newline at end of file
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/test/OpenCL_gpu.c	Sun Feb 05 18:30:30 2017 +0900
@@ -0,0 +1,91 @@
+#include <stdio.h>
+#include <stdlib.h>
+ 
+#ifdef __APPLE__
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#endif
+ 
+#define MEM_SIZE (128)
+#define MAX_SOURCE_SIZE (0x100000)
+ 
+int main()
+{
+    cl_device_id device_id = NULL;
+    cl_context context = NULL;
+    cl_command_queue command_queue = NULL;
+    cl_mem memobj = NULL;
+    cl_program program = NULL;
+    cl_kernel kernel = NULL;
+    cl_platform_id platform_id = NULL;
+    cl_uint ret_num_devices;
+    cl_uint ret_num_platforms;
+    cl_int ret;
+ 
+    char string[MEM_SIZE];
+ 
+    FILE *fp;
+    char fileName[] = "./hello.cl";
+    char *source_str;
+    size_t source_size;
+ 
+/* Load the source code containing the kernel*/
+    fp = fopen(fileName, "r");
+    if (!fp) {
+        fprintf(stderr, "Failed to load kernel.\n");
+        exit(1);
+    }
+    source_str = (char*)malloc(MAX_SOURCE_SIZE);
+    source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
+    fclose(fp);
+ 
+/* Get Platform and Device Info */
+    ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
+    ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
+ 
+/* Create OpenCL context */
+    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
+ 
+/* Create Command Queue */
+    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
+ 
+/* Create Memory Buffer */
+    memobj = clCreateBuffer(context, CL_MEM_READ_WRITE,MEM_SIZE * sizeof(char), NULL, &ret);
+ 
+/* Create Kernel Program from the source */
+    program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
+                                        (const size_t *)&source_size, &ret);
+ 
+/* Build Kernel Program */
+    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
+ 
+/* Create OpenCL Kernel */
+    kernel = clCreateKernel(program, "hello", &ret);
+ 
+/* Set OpenCL Kernel Parameters */
+    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
+ 
+/* Execute OpenCL Kernel */
+    ret = clEnqueueTask(command_queue, kernel, 0, NULL,NULL);
+ 
+/* Copy results from the memory buffer */
+    ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0,
+                              MEM_SIZE * sizeof(char),string, 0, NULL, NULL);
+ 
+/* Display Result */
+    puts(string);
+ 
+/* Finalization */
+    ret = clFlush(command_queue);
+    ret = clFinish(command_queue);
+    ret = clReleaseKernel(kernel);
+    ret = clReleaseProgram(program);
+    ret = clReleaseMemObject(memobj);
+    ret = clReleaseCommandQueue(command_queue);
+    ret = clReleaseContext(context);
+ 
+    free(source_str);
+ 
+    return 0;
+}