Archive
Assignment 3 (CS427) – Game of Life
1. โปรแกรม Game of Life ที่เป็นการทำงานแบบ Sequential จำเป็นจะต้องมีอาร์เรย์จำนวน 2 ชุดในการเก็บข้อมูลของเวลาที่ต่างกันในการทำงาน และใช้ Pointer ในการสลับอาร์เรย์เพื่อประมวลผล
#include <stdio.h>
#include <time.h>
#define MAX_WIDTH 1000
FILE *f;
char board1[MAX_WIDTH * MAX_WIDTH];
char board2[MAX_WIDTH * MAX_WIDTH];
char *currentBoard;
char *pastBoard;
int maxx, maxy;
void init(); /* initial board */
void rand_init_pop(int rand_size); /* initial population */
int in_bound(int i, int j);
int neighbor_num(int i, int j);
void process();
void print_board();
void write_to_file();
int main(){
int i, t = 10;
int randnum = 0;
printf("Enter board width : ");
scanf("%d", &maxx);
printf("Enter board height : ");
scanf("%d", &maxy);
printf("Enter number of random population : ");
scanf("%d", &randnum);
printf("Enter time step : ");
scanf("%d", &t);
pastBoard = board1;
currentBoard = board2;
init();
rand_init_pop(randnum);
f = fopen("sequential_board.txt", "w");
for(i=0;i<t;i++){
fprintf(f, "t = %d\n", i);
write_to_file();
process();
}
fprintf(f, "t = %d\n", i);
write_to_file();
fclose(f);
return 0;
}
void init(){
int i, j;
for(i=0;i<maxy;i++) for(j=0;j<maxx;j++)
board1[i * maxx + j] = board2[i * maxx + j] = 'O';
}
void rand_init_pop(int rand_size){
int i, x, y;
srand( time(0));
for(i=0;i<rand_size;i++){
y = rand() % maxy;
x = rand() % maxx;
if(pastBoard[y * maxx + x] == 'X') i--;
else pastBoard[y * maxx + x] = 'X';
}
}
int in_bound(int i, int j){
if(i<0 || j=maxy || j>=maxy) return 0;
else return 1;
}
int neighbor_num(int i, int j){
int num = 0;
if( in_bound(i-1,j-1) && pastBoard[(i-1) * maxx + (j-1)]=='X' ) num++;
if( in_bound(i-1, j ) && pastBoard[(i-1) * maxx + ( j )]=='X' ) num++;
if( in_bound( i ,j-1) && pastBoard[( i ) * maxx + (j-1)]=='X' ) num++;
if( in_bound(i+1,j+1) && pastBoard[(i+1) * maxx + (j+1)]=='X' ) num++;
if( in_bound(i+1, j ) && pastBoard[(i+1) * maxx + ( j )]=='X' ) num++;
if( in_bound( i ,j+1) && pastBoard[( i ) * maxx + (j+1)]=='X' ) num++;
if( in_bound(i+1,j-1) && pastBoard[(i+1) * maxx + (j-1)]=='X' ) num++;
if( in_bound(i-1,j+1) && pastBoard[(i-1) * maxx + (j+1)]=='X' ) num++;
return num;
}
void process(){
int i,j;
char* tmp;
for(i=0;i<maxy;i++){
for(j=0;j<maxx;j++){
//For a space that is 'populated'
if(pastBoard[i * maxx + j]=='X'){
//Each cell with one or no neighbors dies, as if by loneliness.
if(neighbor_num(i,j) <= 1) currentBoard[i * maxx + j] = 'O';
//Each cell with two or three neighbors survives.
else if(neighbor_num(i,j) <= 3) currentBoard[i * maxx + j] = 'X';
// Each cell with four or more neighbors dies, as if by overpopulation.
else currentBoard[i * maxx + j] = 'O';
}
//For a space that is 'empty' or 'unpopulated'
else{
// Each cell with three neighbors becomes populated.
if(neighbor_num(i,j) == 3) currentBoard[i * maxx + j] = 'X';
else currentBoard[i * maxx + j] = 'O';
}
}
}
tmp = pastBoard;
pastBoard = currentBoard;
currentBoard = tmp;
}
void write_to_file(){
int i, j;
for(i=0;i<maxy;i++){
for(j=0;j<maxx;j++){
if(pastBoard[i * maxx + j]=='X') fprintf(f, "1”);
else fprintf(f, "0”);
}
fprintf(f, "\n");
}
fprintf(f, "\n");
}
เมื่อเริ่มต้นโปรแกรม โปรแกรมจะให้ใส่ข้อมูลขนาดความกว้าง และความสูงของตารางที่จะใช้ จากนั้นจะถามถึงจำนวน Cell ที่มีชีวิตเริ่มต้นที่ต้องการกำหนด โดยจะมีตำแหน่งแบบสุ่ม และสุดท้ายจะถามถึง จำนวนหน่วยเวลาที่ต้องการทำการ Simulation เมื่อใส่ข้อมูลครบแล้วโปรแกรมจะทำงาน และเขียนข้อมูลของตารางในแต่ละช่วงเวลาไปยังไฟล์ที่ชื่อ “sequential_board.txt”
2. การทำงานของโปรแกรม Game of Life บนจีพียู ใน 1 time step จะต้องเข้าถึงหน่วยความจำเฉลี่ยประมาณ 8 ครั้งต่อ 1 เทรด ดังนั้นเพื่อให้มีความรวดเร็วในการทำงาน จะต้องทำการคัดลอดข้อมูลจาก Global Memory มายัง Shared Memory ก่อน เพื่อให้เข้าถึงหน่วยความจำได้เร็วขึ้น โดยวิธีที่ใช้คือการแบ่งการทำงานออกเป็นเทรดบล็อคโดยที่ในตอนเริ่มต้นการทำงาน แต่ละเทรดบล็อคจะต้องคัดลอกข้อมูลในส่วนของ cell ที่มีพื้นที่ติดกันกับพื้นที่ที่เทรดบล็อคนั้นรับผิดชอบมาด้วย แต่ในการเขียนข้อมูลจะเขียนข้อมูลลงใน Global Memory ในช่องที่เทรดบล็อคนั้นรับผิดชอบเท่านั้น
#include <stdio.h>
#include <math.h>
#include <time.h>
#define MAX_WIDTH 500
#define TILE_WIDTH 16
#define TY threadIdx.y
#define TX threadIdx.x
void rand_init_pop(int *board, int maxx, int maxy, int rand_size);
void write_to_file(FILE *f, int *board, int maxx, int maxy);
__global__ void process(int *board, int maxx, int maxy);
int main(){
int *board, *d_board, i, t;
int maxx, maxy, size, randnum;
FILE *f;
printf("Enter board width : ");
scanf("%d", &maxx);
printf("Enter board height : ");
scanf("%d", &maxy);
printf("Enter number of random population : ");
scanf("%d", &randnum);
printf("Enter time step : ");
scanf("%d", &t);
size = sizeof(int) * maxx * maxy;
f = fopen("cuda_board.txt", "w");
board = (int*)malloc(size);
cudaMalloc((void**)&d_board, size);
rand_init_pop(board, maxx, maxy, randnum);
dim3 blockDim(TILE_WIDTH+2, TILE_WIDTH+2);
dim3 gridDim(ceil(maxy*1.0/TILE_WIDTH), ceil(maxx*1.0/TILE_WIDTH));
cudaMemcpy(d_board, board, size, cudaMemcpyHostToDevice);
for(i=0;i<t;i++){
fprintf(f, "t = %d\n", i);
write_to_file(f, board, maxx, maxy);
process<<>>(d_board, maxx, maxy);
cudaMemcpy(board, d_board, size, cudaMemcpyDeviceToHost);
}
fprintf(f, "t = %d\n", i);
write_to_file(f, board, maxx, maxy);
fclose(f);
free(board);
cudaFree(d_board);
return 0;
}
void init(int *board, int maxx, int maxy){
int i, j;
for(i=0;i<maxy;i++) for(j=0;j<maxx;j++)
board[i * maxx + j] = 0;
}
void rand_init_pop(int *board, int maxx, int maxy, int rand_size){
int i, x, y;
srand( time(0));
init(board, maxx, maxy);
for(i=0;i<rand_size;i++){
y = rand() % maxy;
x = rand() % maxx;
if(board[y * maxx + x] == 1) i--;
else board[y * maxx + x] = 1;
}
}
void write_to_file(FILE *f, int *board, int maxx, int maxy){
for(int i=0;i<maxy;i++){
for(int j=0;j<maxx;j++)
fprintf(f, "%d", board[i * maxx + j]);
fprintf(f, "\n");
}
fprintf(f, "\n");
}
__global__ void process(int *board, int maxx, int maxy){
__shared__ int sboard[TILE_WIDTH+2][TILE_WIDTH+2];
int i = (blockIdx.x * TILE_WIDTH) + TX - 1;
int j = (blockIdx.y * TILE_WIDTH) + TY - 1;
if(i<=maxy && j=0 && j>=0 && i<maxy && j0 && TX0 && TY<TILE_WIDTH){
int n = 0;
n += sboard[TY+1][TX+1];
n += sboard[TY+1][TX-1];
n += sboard[TY-1][TX-1];
n += sboard[TY-1][TX+1];
n += sboard[TY+1][ TX ];
n += sboard[ TY ][TX+1];
n += sboard[ TY ][TX-1];
n += sboard[TY-1][ TX ];
if(sboard[TY][TX]==1 && n!=2 && n!=3)
board[boardidx] = 0;
else if(sboard[TY][TX]==0 && n==3)
board[boardidx] = 1;
}
}
}
เมื่อรันโปรแกรม โปรแกรมจะถามข้อมูลเช่นเดียวกับโปรแกรม Sequential แต่ข้อมูลจะเขียนลงไปที่ไฟล์ที่มีชื่อว่า “cuda_board.txt”
สามารถดาวน์โหลดไฟล์โปรแกรมได้ที่ด้านล่าง
GameOfLife.c
GameOfLifeCUDA.cu
Assignment 2 (CS427)
1. การทำให้โปรแกรมมีลักษณะการทำงานแบบ Parallelism มากที่สุด จะต้องทำการวิเคราะห์ความสัมพันธ์ของโปรแกรม ในส่วนต่างๆว่ามีส่วนใดที่สามารถทำงานพร้อมกันได้โดยดูจากกราฟ
จากโปรแกรมตัวอย่างเมื่อเขียนความสัมพันธ์ในรูปแบบกราฟจะได้กราฟลักษณะข้างต้น ซึ่งหมายความว่าลูปที่ 2 และ 3 ที่เป็นส่วนที่คำนวนค่าของ c และ d สามารถทำงานพร้อมกันได้ ดังนั้นจึงสามารถใช้คำสั่ง #pragma omp sections ในส่วนนี้ และภายในลูปแต่ละลูปก็มีไม่ความสัมพันธ์กันระหว่างแต่ละการวนรอบสังเกตุได้จากการอ้างอิงถึง index เดียวกันในแต่ละการวนรอบเท่านั้น ดังนั้นทุกลูปสามารถใช้คำสั่ง #pragma omp for เพื่อให้เกิดการทำงานแบบขนานขึ้นในแต่ละลูป
โดยสามารถเขียนเป็นโปรแกรมได้ดังดังนี้
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>
#define N 50
int main (int argc, char *argv[])
{
int i, nthreads, tid;
float a[N], b[N], c[N], d[N];
float pa[N], pb[N], pc[N], pd[N];
/* Sequential Program */
for (i=0; i<N; i++) {
a[i] = i * 1.5;
b[i] = i + 22.35;
c[i] = d[i] = 0.0;
}
for (i=0; i<N; i++){
c[i] = a[i] + b[i];
}
for (i=0; i<N; i++){
d[i] = a[i] * b[i];
}
/* Parallel Program with OpenMP */
#pragma omp parallel for
for (i=0; i<N; i++) {
pa[i] = i * 1.5;
pb[i] = i + 22.35;
pc[i] = pd[i] = 0.0;
}
#pragma omp parallel sections
{
#pragma omp section
#pragma omp parallel for
for (i=0; i<N; i++){
pc[i] = pa[i] + pb[i];
}
#pragma omp section
#pragma omp parallel for
for (i=0; i<N; i++){
pd[i] = pa[i] * pb[i];
}
}
/* check answers */
for(i=0;i<N;i++){
if(c[i]!=pc[i] || d[i]!=pd[i]){
printf("Compute Error\n");
break;
}
}
if(i==N) printf("Corrected Answer\n");
return 0;
}
โปรแกรมดังกล่าวจะทำงานทั้งหมดสองครั้งคือ คำนวนเชิงลำดับ (Sequential) และคำนวนแบบขนาน (Parallel) ด้วยการใช้ OpenMP หลังจากนั้นจะนำคำตอบมาตรวจสอบว่าการคำนวนแบบขนานให้ผลลัพท์เช่นเดียวกับการทำงานเชิงลำดับหรือไม่ ซึ่งหากคำนวนได้ถูกต้องจะปรากฎผลลัพท์ขึ้นมาเป็น Corrected Answer และจะปรากฎผลลัพท์ว่า Compute Error หากทำงานผิดพลาด ในที่นี้จะทำการคอมไพล์ด้วยโปรแกรม gcc เวอร์ชัน 4 ด้วยคำสั่ง
$gcc-4 –fopenmp program1.c –o program1
และเมื่อนำโปรแกรมไปรันทดสอบ ผลลัพท์ที่ได้คือ Corrected Answer หรือโปรแกรมทำงานได้ถูกต้อง ถึงแม้โปรแกรมนี้จะมีความเป็น Parallelism สูง แต่ในการทำงานจริงจะทำงานได้ช้า เนื่องจากโปรแกรมจะเสียเวลาในการสร้างเทรดขึ้นมาจำนวนมาก ซึ่งไม่คุ้มกับการคำนวนข้อมูลที่มีจำนวนน้อย ในการทำงานจริงเพื่อให้มีประสิทธิภาพสูงสุดอาจจะกำหนดให้มีการทำงานแบบขนานเพียงบางส่วนเท่านั้น เพื่อลด overhead ในการสร้างเทรด
2. ในข้อนี้จะทำการเปลี่ยนโปรแกรมเชิงลำดับให้เป็นโปรแกรมแบบขนานที่ทำงานได้บนจีพียู ซึ่งจะเขียนโปรแกรมที่คำนวนบนซีพียูและจีพียูจากนั้นจะนำคำตอบที่ได้มาตรวจสอบว่าได้ค่าที่เหมือนกัน ซึ่งการคำนวนค่าของ Floating-point บนซีพียูและจีพียูนั้นให้ค่าที่แตกต่างกันเล็กน้อย ดังนั้นเทียบค่าของ Floating-point ที่คำนวนได้ด้วยฟังก์ชัน floatIsEqual ในโปรแกรม
#include <stdio.h>
#include <time.h>
#define DELTA 0.001
int floatIsEqual(float a, float b){
if( (b<a+DELTA) && (b>a-DELTA) ) return 1;
else return 0;
}
__global__ void GPUfunction(float *A, float *B, float *C){
int i = blockIdx.x * blockDim.x + threadIdx.x;
if(i<1024) C[i] = A[i] + B[i];
else if(i<2048) C[i] = A[i] * B[i];
else if(i<3072) C[i] = A[i] - B[i];
else if(i<4096) C[i] = A[i] / B[i];
}
void CPUfunction(float *A, float *B, float *C){
int i;
for(i=0; i<1024; i++) C[i] = A[i] + B[i];
for(i=1024; i<2048; i++) C[i] = A[i] * B[i];
for(i=2048; i<3072; i++) C[i] = A[i] - B[i];
for(i=3072; i<4096; i++) C[i] = A[i] / B[i];
}
int main(){
int i;
float A[4096], B[4096], C[4096];
float *hA, *hB, *hC;
float *dA, *dB, *dC;
int size = 4096 * sizeof(float);
hA = (float*)malloc(size);
hB = (float*)malloc(size);
hC = (float*)malloc(size);
cudaMalloc((void**)&dA, size);
cudaMalloc((void**)&dB, size);
cudaMalloc((void**)&dC, size);
/* initialize */
srand(time(0));
for(i=0;i<4096;i++){
A[i] = hA[i] = rand() % 40;
// make sure that B is not zero
B[i] = hB[i] = rand() % 39 + 1;
}
/* GPU code */
cudaMemcpy(dA, hA, size, cudaMemcpyHostToDevice);
cudaMemcpy(dB, hB, size, cudaMemcpyHostToDevice);
GPUfunction<<<8, 512>>>(dA, dB, dC);
cudaMemcpy(hC, dC, size, cudaMemcpyDeviceToHost);
/* CPU code */
CPUfunction(A, B, C);
/* check answers from CPU and GPU are same */
for(i=0;i<4096;i++){
if(!floatIsEqual(C[i], hC[i])){
printf("Compute Error\n");
break;
}
}
if(i==4096) printf("Corrected Answer\n");
free(hA);
free(hB);
free(hC);
cudaFree(dA);
cudaFree(dB);
cudaFree(dC);
return 0;
}
ทำการคอมไพล์โปรแกรมด้วยคำสั่ง
$nvcc program2.cu –o program2
ได้ผลลัพท์ออกมาเป็น Corrected Answer ซึ่งหมายถึงโปรแกรมทำงานได้ถูกต้อง การคำนวนแบบขนานจากจีพียูให้ผลลัพท์ตรงตามการคำนวนจากซีพียู
3. โปรแกรมนี้สามารถเพิ่มประสิทธิภาพได้ด้วยการใช้ Shared Memory และ Parallel Reduction ในการรวมข้อมูลหลังจากนั้นจึงทำการคำนวน โค้ดของโปรแกรมเหมือนกับข้อ 2 ยกเว้นในส่วนของฟังก์ชัน GPUfunction และ CPUfunction ซึ่งทำการเปลี่ยนแปลงดังนี้
GPUfunction :
__global__ void GPUfunction(float *A, float *B, float *C){
int i = blockIdx.x * blockDim.x + threadIdx.x;
__shared__ float ssuma[512];
float sum;
int s, tid = threadIdx.x;
/* copy global mem to shared mem synchonously */
ssuma[tid] = A[i];
__syncthreads();
/* parallel reduction */
for(s=256; s>0; s/=2){
if(tid<s) ssuma[tid] = ssuma[tid] + ssuma[tid+s];
__syncthreads();
}
sum = ssuma[0];
if(i<1024) C[i] = sum + B[i];
else if(i<2048) C[i] = sum * B[i];
else if(i<3072) C[i] = sum - B[i];
else if(i<4096) C[i] = sum / B[i];
}
CPUfunction :
void CPUfunction(float *A, float *B, float *C){
int i;
float SUMA[8];
/* initialize SUMA */
for(i=0;i<8;i++) SUMA[i] = 0.0;
for(i=0;i<4096;i++) SUMA[i/512] += A[i];
for(i=0; i<1024; i++)
C[i] = SUMA[i/512] + B[i];
for(i=1024; i<2048; i++)
C[i] = SUMA[i/512] * B[i];
for(i=2048; i<3072; i++)
C[i] = SUMA[i/512] - B[i];
for(i=3072; i<4096; i++)
C[i] = SUMA[i/512] / B[i];
}
ทำการคอมไพล์โปรแกรมด้วย nvcc โดยใช้คำสั่ง
$nvcc program3.cu –o program3
ซึ่งเมื่อรันโปรแกรมแล้วให้ผลลัพท์ถูกต้อง และแสดงข้อความว่า Corrected Answer เช่นเดียวกันในข้อ 2
ดาวน์โหลดโปรแกรมทั้ง 3 โปรแกรมได้ที่ด้านล่าง
program 1
program 2
program 3
ขั้นตอนการติดตั้ง CUDA Toolkit และ NVIDIA CUDA SDK
1. หลังจากที่ Log in เข้าไปที่เครื่องเซิฟเวอร์แล้ว ทำการสร้างไดเรคเทอรี่ส่วนตัวขึ้นมา (ในที่นี้จะเก็บไว้ที่ /tmp/4909611727) ทำการคัดลอกไฟล์ cudatoolkit_2.3_linux_64_ubuntu9.04.run และ cudasdk_2.3_linux.run จาก /home/kasidit มาไว้ยังไดเรคเทอรี่ที่สร้างขึ้น
2. ติดตั้ง CUDA Toolkit ด้วยคำสั่ง sh cudatoolkit_2.3_linux_64_ubuntu9.04.run จากนั้นโปรแกรมจะถามถึง path ที่ต้องการจะใช้ติดตั้ง พิมพ์ /tmp/4909611727 ลงไป
3. คอมไพเลอร์ของ CUDA หรือ nvcc นั้นจะเก็บไว้ที่ /tmp/4909611727/cuda/bin ให้ตั้งค่าของ $PATH และ $LD_LIBRARY_PATH ดังนี้
export PATH=/tmp/4909611727/cuda/bin:$PATH
export LD_LIBRARY_PATH=/tmp/4909611727/cuda/lib64:$LD_LIBRARY_PATH
4. ทดสอบว่าสามารถใช้งาน nvcc ได้ด้วยการพิมพ์คำสั่ง nvcc –version ซึ่งจะได้ผลลัพท์ดังนี้
5. ติดตั้ง NVIDIA CUDA SDK ด้วยคำสั่ง sh cudasdk_2.3_linux.run โปรแกรมจะถามถึง path ที่จะใช้ติดตั้ง ซึ่งถ้าไม่ใส่จะถูกกำหนดให้เป็น NVIDIA_GPU_Computing_SDK แต่ในที่นี้จะกำหนดให้เป็น /tmp/4909611727/NVIDAI_CUDA_SDK
6. คอมไพล์ไฟล์ตัวอย่างที่มากับ NVIDIA CUDA SDK ด้วยการเข้าไปที่ /tmp/4909611727/NVIDIA_CUDA_SDK/C แล้วใช้คำสั่ง make
7. ทดสอบโปรแกรมด้วยการรันโปรแกรมตัวอย่างที่อยู่ที่ /tmp/4909611727/NVIDIA_CUDA_SDK/C/bin/linux/release ซึ่งโปรแกรมตัวอย่างที่จะทดสอบคือ deviceQuery และ bandwidthTest
deviceQuery
โปรแกรม deviceQuery เป็นโปรแกรมที่ใช้แสดงคุณลักษณะของกราฟิกการ์ดที่ใช้ ซึ่งผลลัพธ์จากการรันโปรแกรม deviceQuery เป็นดังต่อไปนี้
bandwidthTest
โปรแกรม bandwidthTest จะแสดงให้เห็นถึงความสามารถในการส่งข้อมูลแบบต่างๆระหว่างหน่วยความจำหลัก และหน่วยความจำ Global ของ GPU ทดสอบโดยการส่งข้อมูลขนาด 32 MB ผลได้เป็นดังโปรแกรม
Vector Addition
เขียนโปรแกรมบวกเวกเตอร์ เซฟด้วยชื่อไฟล์ vecadd.cu แล้วทำการคัดลอกไปยังเครื่องเซิฟเวอร์ และคอมไพล์ด้วยคำสั่ง nvcc vecadd.cu และรันโปรแกรมจะได้ผลลัพท์ออกมาดังนี้
โค้ดของโปรแกรมนี้คือ
#include <stdio.h>
__global__ void VecAdd(float *A, float *B, float *C, int N){
int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
if(idx<N){
C[idx] = A[idx] + B[idx];
}
}
int main(){
float *h_A, *h_B, *h_C;
float *d_A, *d_B, *d_C;
int N = 200000;
int size = sizeof(float) * N;
int threadPerBlock = 512;
int blockPerGrid = (N + threadPerBlock -1) / threadPerBlock;
// Allocate host and device memory
h_A = (float*)malloc(size); h_B = (float*)malloc(size); h_C = (float*)malloc(size);
cudaMalloc((void**)&d_A, size); cudaMalloc((void**)&d_B, size); cudaMalloc((void**)&d_C, size);
// Initial value on vector A and B
for(int i=0;i<N;i++)
h_A[i] = rand()/200, h_B[i] = rand()/200;
// Copy vector A and B to device
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Call kernel
VecAdd<<< blockPerGrid, threadPerBlock >>>(d_A, d_B, d_C, N);
// Copy back vector C
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
printf("%s\n", cudaGetErrorString( cudaGetLastError() ) );
printf("Threads per block = %d\nBlocks per grid = %d\n", threadPerBlock, blockPerGrid);
printf("Vector size = %d\nMemory used = %f MB\n", N, 3 * size / 1024.0 / 1024);
// Deallocate host and device memory
free(h_A); free(h_B); free(h_C);
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
return 0;
}
Remixman 