-
Notifications
You must be signed in to change notification settings - Fork 0
/
my-example-3a.cpp
91 lines (55 loc) · 1.78 KB
/
my-example-3a.cpp
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
/*
please note that the series of optmiztion technology is not in official document.
All the tests are based on AMD MI25 radeon instict and AMD ROCm.
*/
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <hip/hip_runtime.h>
#define HIP_ASSERT(x) (assert((x)==hipSuccess))
#define NUM 1
#define THREADS_PER_BLOCK_X 256
#define THREADS_PER_BLOCK_Y 1
#define THREADS_PER_BLOCK_Z 1
#define MAX_VGPRS 255
__global__ void ljx_test_kernel_255()
{
asm volatile("v_mov_b32 v0, 0");
asm volatile("v_mov_b32 v255, 0" );
}
__global__ void ljx_test_kernel_256()
{
asm volatile("v_mov_b32 v0, 0");
asm volatile("v_mov_b32 v256, 0");
}
using namespace std;
int main() {
float* hostA;
float* deviceA;
// hipDeviceProp_t devProp;
// hipGetDeviceProperties(&devProp, 0);
// cout << " System minor " << devProp.minor << endl;
// cout << " System major " << devProp.major << endl;
// cout << " agent prop name " << devProp.name << endl;
// cout << "hip Device prop succeeded " << endl ;
// hipEvent_t start, stop;
// hipEventCreate(&start);
// hipEventCreate(&stop);
// float eventMs = 1.0f;
// int i;
int errors;
//hostA = (float*)malloc(NUM * sizeof(float));
//HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float)));
int vgprs;
vgprs = 255;
printf(" Begin lauch %d VGPRs passed \n", vgprs);
// my method
dim3 blocksize(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y);
dim3 gridsize(1, 1);
// ljx_test_kernel_255<<<gridsize, blocksize >>> ();
ljx_test_kernel_256<<<gridsize, blocksize >>> ();
printf("%d VGPRs passed \n", vgprs);
return errors;
}