Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

ocl: minor fixes and adjustments #878

Merged
merged 1 commit into from
Jan 6, 2025
Merged

ocl: minor fixes and adjustments #878

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 12 additions & 6 deletions src/acc/opencl/Makefile
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,10 +22,12 @@ endif
UNAME := $(shell uname)
HEADERONLY ?= 0
STATIC ?= 1
INTEL ?= 0
GNU ?= 0
DEV ?= 0

# Intel Compiler
ICX := $(shell which icx 2>/dev/null)
INTEL ?= $(if $(ICX),$(if $(filter-out 0,$(GNU)),0,2),0)

# select from set of predefined triplet specifications
SPECID ?= 0
# limit shape in tests (zero or negative for unlimited)
Expand Down Expand Up @@ -200,11 +202,15 @@ else
endif
else ifeq (,$(OPENCL_INC))
ifneq (,$(wildcard $(OPENCL_ROOT)/include/CL/cl.h))
LDFLAGS += -L$(OPENCL_ROOT)/$(if $(wildcard $(OPENCL_ROOT)/lib64),lib64,lib)
OPENCL_INC := $(OPENCL_ROOT)/include
LDFLAGS += -L$(OPENCL_ROOT)/lib64
else ifneq (,$(wildcard $(OPENCL_ROOT)/include/sycl/CL/cl.h))
OPENCL_INC := $(OPENCL_ROOT)/include/sycl
LDFLAGS += -L$(OPENCL_ROOT)/compiler/lib/intel64 -lintlc
else ifneq (,$(ICX))
OPENCL_ROOT := $(abspath $(dir $(ICX))/..)
ifneq (,$(wildcard $(OPENCL_ROOT)/include/sycl/CL/cl.h))
LDFLAGS += -L$(OPENCL_ROOT)/$(if $(wildcard $(OPENCL_ROOT)/lib64),lib64,lib)
LDFLAGS += -L$(OPENCL_ROOT)/compiler/lib/intel64 -lintlc
OPENCL_INC := $(OPENCL_ROOT)/include/sycl
endif
endif
endif
# OPENCL_INC: directory containing CL/cl.h.
Expand Down
267 changes: 149 additions & 118 deletions src/acc/opencl/acc_opencl.c
View file
Open in desktop

Large diffs are not rendered by default.

4 changes: 4 additions & 0 deletions src/acc/opencl/acc_opencl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -338,8 +338,12 @@ typedef struct c_dbcsr_acc_opencl_config_t {
c_dbcsr_acc_opencl_timer_t timer; /* c_dbcsr_acc_opencl_device_t? */
/** Kernel-parameters are matched against device's UID */
cl_uint devmatch;
/** Split devices into sub-devices (if possible) */
cl_int devsplit;
/** Verbosity level (output on stderr). */
cl_int verbosity;
/** Guessed number of ranks per node (local), and rank-ID. */
cl_int nranks, nrank;
/** Non-zero if library is initialized (negative: no device). */
cl_int ndevices;
/** Maximum number of threads (omp_get_max_threads). */
Expand Down
141 changes: 75 additions & 66 deletions src/acc/opencl/acc_opencl_mem.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -189,7 +189,7 @@ int c_dbcsr_acc_host_mem_allocate(void** host_mem, size_t nbytes, void* stream)
}
# endif
# if defined(ACC_OPENCL_XHINTS)
if (0 != (8 & c_dbcsr_acc_opencl_config.xhints) && (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS))) {
if (0 != (16 & c_dbcsr_acc_opencl_config.xhints) && (0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS))) {
host_ptr = malloc(nbytes);
if (NULL != host_ptr) flags = CL_MEM_USE_HOST_PTR;
}
Expand All @@ -202,7 +202,7 @@ int c_dbcsr_acc_host_mem_allocate(void** host_mem, size_t nbytes, void* stream)
: c_dbcsr_acc_opencl_stream_default());
mapped = clEnqueueMapBuffer(str->queue, memory, CL_TRUE /*always block*/,
# if defined(ACC_OPENCL_XHINTS) && (defined(CL_VERSION_1_2) || defined(CL_MAP_WRITE_INVALIDATE_REGION))
(4 & c_dbcsr_acc_opencl_config.xhints) ? CL_MAP_WRITE_INVALIDATE_REGION :
(8 & c_dbcsr_acc_opencl_config.xhints) ? CL_MAP_WRITE_INVALIDATE_REGION :
# endif
(CL_MAP_READ | CL_MAP_WRITE),
0 /*offset*/, nbytes, 0, NULL, NULL, &result);
Expand Down Expand Up @@ -250,7 +250,7 @@ int c_dbcsr_acc_host_mem_deallocate(void* host_mem, void* stream) {
void* host_ptr = NULL;
int result_release;
# if defined(ACC_OPENCL_XHINTS)
if (0 != (8 & c_dbcsr_acc_opencl_config.xhints) &&
if (0 != (16 & c_dbcsr_acc_opencl_config.xhints) &&
(0 != c_dbcsr_acc_opencl_config.device.nv || NULL != (ACC_OPENCL_XHINTS)) &&
EXIT_SUCCESS == clGetMemObjectInfo(info.memory, CL_MEM_HOST_PTR, sizeof(void*), &host_ptr, NULL) && NULL != host_ptr)
{
Expand Down Expand Up @@ -343,79 +343,88 @@ int c_dbcsr_acc_dev_mem_allocate(void** dev_mem, size_t nbytes) {
}
# endif
assert(NULL != dev_mem && NULL != context);
if (0 != nbytes) {
# if defined(ACC_OPENCL_MEM_DEVPTR)
if (NULL != c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL) {
assert(0 == c_dbcsr_acc_opencl_config.device.unified);
*dev_mem = memptr = c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL(
context, c_dbcsr_acc_opencl_config.device.id, NULL /*properties*/, nbytes, 0 /*alignment*/, &result);
if (EXIT_SUCCESS != result) *dev_mem = NULL;
}
else
# endif
{
const int devuid = c_dbcsr_acc_opencl_config.device.uid;
const int try_flag = ((0 != c_dbcsr_acc_opencl_config.device.unified || 0 == c_dbcsr_acc_opencl_config.device.intel ||
(0x4905 != devuid && 0x020a != devuid && (0x0bd0 > devuid || 0x0bdb < devuid)))
? 0
: (1u << 22));
memory = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE | try_flag), nbytes, NULL /*host_ptr*/, &result);
if (0 != try_flag && EXIT_SUCCESS != result) { /* retry without try_flag */
memory = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL /*host_ptr*/, &result);
if (NULL != c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL) {
assert(0 == c_dbcsr_acc_opencl_config.device.unified);
*dev_mem = memptr = c_dbcsr_acc_opencl_config.device.clDeviceMemAllocINTEL(
context, c_dbcsr_acc_opencl_config.device.id, NULL /*properties*/, nbytes, 0 /*alignment*/, &result);
if (EXIT_SUCCESS != result) *dev_mem = NULL;
}
if (EXIT_SUCCESS == result) {
# if defined(ACC_OPENCL_MEM_DEVPTR)
const c_dbcsr_acc_opencl_stream_t* str = NULL;
static cl_kernel kernel = NULL;
const size_t size = 1;
ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_memory);
str = c_dbcsr_acc_opencl_stream(NULL /*lock*/, ACC_OPENCL_OMP_TID());
/* determine device-side value of device-memory object by running some kernel */
assert(NULL != memory && NULL != str && NULL != str->queue);
if (NULL == kernel) { /* generate kernel */
const char source[] = "kernel void memptr(global unsigned long* ptr) {\n"
" const union { global unsigned long* p; unsigned long u; } cast = { ptr };\n"
" const size_t i = get_global_id(0);\n"
" ptr[i] = cast.u + i;\n"
"}\n";
assert(sizeof(size_t) == sizeof(cl_ulong));
result = c_dbcsr_acc_opencl_kernel(0 /*source_is_file*/, source, "memptr" /*kernel_name*/, NULL /*build_params*/,
NULL /*build_options*/, NULL /*try_build_options*/, NULL /*try_ok*/, NULL /*extnames*/, 0 /*num_exts*/, &kernel);
}
/* TODO: backup/restore memory */
if (EXIT_SUCCESS == result) result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory);
if (EXIT_SUCCESS == result) {
result = clEnqueueNDRangeKernel(
str->queue, kernel, 1 /*work_dim*/, NULL /*offset*/, &size, NULL /*local_work_size*/, 0, NULL, NULL);
}
if (EXIT_SUCCESS == result) {
result = c_dbcsr_acc_opencl_memcpy_d2h(memory, &memptr, 0, sizeof(void*), str->queue, 1 /*blocking*/);
else
# endif
{
const int devuid = c_dbcsr_acc_opencl_config.device.uid;
const int try_flag = ((0 != c_dbcsr_acc_opencl_config.device.unified || 0 == c_dbcsr_acc_opencl_config.device.intel ||
(0x4905 != devuid && 0x020a != devuid && (0x0bd0 > devuid || 0x0bdb < devuid)))
? 0
: (1u << 22));
memory = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE | try_flag), nbytes, NULL /*host_ptr*/, &result);
if (0 != try_flag && EXIT_SUCCESS != result) { /* retry without try_flag */
memory = clCreateBuffer(context, CL_MEM_READ_WRITE, nbytes, NULL /*host_ptr*/, &result);
}
assert(EXIT_SUCCESS != result || NULL != memptr);
if (EXIT_SUCCESS == result) {
c_dbcsr_acc_opencl_info_memptr_t* const info = (c_dbcsr_acc_opencl_info_memptr_t*)c_dbcsr_acc_opencl_pmalloc(
NULL /*lock*/, (void**)c_dbcsr_acc_opencl_config.memptrs, &c_dbcsr_acc_opencl_config.nmemptrs);
assert(NULL != memptr);
if (NULL != info) {
info->memory = memory;
info->memptr = memptr;
*dev_mem = memptr;
# if defined(ACC_OPENCL_MEM_DEVPTR)
const c_dbcsr_acc_opencl_stream_t* str = NULL;
static cl_kernel kernel = NULL;
const size_t size = 1;
ACC_OPENCL_ACQUIRE(c_dbcsr_acc_opencl_config.lock_memory);
str = c_dbcsr_acc_opencl_stream(NULL /*lock*/, ACC_OPENCL_OMP_TID());
/* determine device-side value of device-memory object by running some kernel */
assert(NULL != memory && NULL != str && NULL != str->queue);
if (NULL == kernel) { /* generate kernel */
const char source[] = "kernel void memptr(global unsigned long* ptr) {\n"
" const union { global unsigned long* p; unsigned long u; } cast = { ptr };\n"
" const size_t i = get_global_id(0);\n"
" ptr[i] = cast.u + i;\n"
"}\n";
assert(sizeof(size_t) == sizeof(cl_ulong));
result = c_dbcsr_acc_opencl_kernel(0 /*source_is_file*/, source, "memptr" /*kernel_name*/, NULL /*build_params*/,
NULL /*build_options*/, NULL /*try_build_options*/, NULL /*try_ok*/, NULL /*extnames*/, 0 /*num_exts*/, &kernel);
}
else result = EXIT_FAILURE;
}
ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_memory);
/* TODO: backup/restore memory */
if (EXIT_SUCCESS == result) result = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory);
if (EXIT_SUCCESS == result) {
result = clEnqueueNDRangeKernel(
str->queue, kernel, 1 /*work_dim*/, NULL /*offset*/, &size, NULL /*local_work_size*/, 0, NULL, NULL);
}
if (EXIT_SUCCESS == result) {
result = c_dbcsr_acc_opencl_memcpy_d2h(memory, &memptr, 0, sizeof(void*), str->queue, 1 /*blocking*/);
}
assert(EXIT_SUCCESS != result || NULL != memptr);
if (EXIT_SUCCESS == result) {
c_dbcsr_acc_opencl_info_memptr_t* const info = (c_dbcsr_acc_opencl_info_memptr_t*)c_dbcsr_acc_opencl_pmalloc(
NULL /*lock*/, (void**)c_dbcsr_acc_opencl_config.memptrs, &c_dbcsr_acc_opencl_config.nmemptrs);
assert(NULL != memptr);
if (NULL != info) {
info->memory = memory;
info->memptr = memptr;
*dev_mem = memptr;
}
else result = EXIT_FAILURE;
}
ACC_OPENCL_RELEASE(c_dbcsr_acc_opencl_config.lock_memory);
# else
*dev_mem = memptr = memory;
*dev_mem = memptr = memory;
# endif
}
if (EXIT_SUCCESS != result) {
if (NULL != memory) ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseMemObject(memory));
*dev_mem = NULL;
}
}
if (EXIT_SUCCESS != result) {
if (NULL != memory) ACC_OPENCL_EXPECT(EXIT_SUCCESS == clReleaseMemObject(memory));
*dev_mem = NULL;
if (0 != c_dbcsr_acc_opencl_config.verbosity) {
if (EXIT_SUCCESS == result && 0 != c_dbcsr_acc_opencl_config.debug) {
fprintf(stderr, "INFO ACC/OpenCL: memory=%p pointer=%p size=%llu successfully allocated\n", (const void*)memory, memptr,
(unsigned long long)nbytes);
}
else if (EXIT_SUCCESS != result) {
fprintf(stderr, "ERROR ACC/OpenCL: memory=%p pointer=%p size=%llu failed to allocate\n", (const void*)memory, memptr,
(unsigned long long)nbytes);
}
}
}
if (0 != c_dbcsr_acc_opencl_config.debug && 0 != c_dbcsr_acc_opencl_config.verbosity && EXIT_SUCCESS == result) {
fprintf(stderr, "INFO ACC/OpenCL: memory=%p pointer=%p size=%llu allocated\n", (const void*)memory, memptr,
(unsigned long long)nbytes);
}
else *dev_mem = NULL;
assert(EXIT_SUCCESS == result || NULL == *dev_mem);
# if defined(__DBCSR_ACC) && defined(ACC_OPENCL_PROFILE)
c_dbcsr_timestop(&routine_handle);
Expand Down
2 changes: 1 addition & 1 deletion src/acc/opencl/acc_opencl_stream.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -118,7 +118,7 @@ int c_dbcsr_acc_stream_create(void** stream_p, const char* name, int priority) {
# endif
{
# if defined(ACC_OPENCL_XHINTS)
if ((2 & c_dbcsr_acc_opencl_config.xhints) && 0 != c_dbcsr_acc_opencl_config.device.intel) { /* enable queue families */
if ((4 & c_dbcsr_acc_opencl_config.xhints) && 0 != c_dbcsr_acc_opencl_config.device.intel) { /* enable queue families */
struct {
cl_command_queue_properties properties;
cl_bitfield capabilities;
Expand Down
24 changes: 13 additions & 11 deletions src/acc/opencl/smm/opencl_libsmm.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -451,20 +451,24 @@ int libsmm_acc_init(void) {
}
# if defined(OPENCL_KERNELS_PARAMS_SMM) && defined(OPENCL_KERNELS_DEVICES)
if (EXIT_SUCCESS == result && (0 == ntuned || 0 != key_direct_skip)) {
unsigned int default_uid = c_dbcsr_acc_opencl_config.device.uid;
const char *line = OPENCL_KERNELS_PARAMS_SMM, *next;
# if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
int active_match = -1;
if (EXIT_SUCCESS == c_dbcsr_acc_opencl_device_name(c_dbcsr_acc_opencl_config.device.id, bufname, ACC_OPENCL_BUFFERSIZE,
NULL /*platform*/, 0 /*platform_maxlen*/, /*cleanup*/ 1))
{ /* determine best-matching parameters based on name of device */
int i = 0, best = 0;
if (1 >= c_dbcsr_acc_opencl_config.devmatch) {
c_dbcsr_acc_opencl_device_uid(c_dbcsr_acc_opencl_config.device.id, bufname, &default_uid);
}
for (; i < ndevices_params; ++i) {
const int score = libxsmm_strimatch(bufname, OPENCL_KERNELS_DEVICES[i], NULL);
unsigned int uid;
if (best < score ||
((best == score) &&
EXIT_SUCCESS == c_dbcsr_acc_opencl_device_uid(NULL /*device*/, OPENCL_KERNELS_DEVICES[i], &uid) &&
uid == c_dbcsr_acc_opencl_config.device.uid))
uid == default_uid))
{
active_match = i;
best = score;
Expand All @@ -489,22 +493,20 @@ int libsmm_acc_init(void) {
{
key.devuid = 0;
}
config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key));
config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key)); /* duplicate? */
if (NULL == config_init) {
if (NULL == libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) {
if (NULL != libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) ++ntuned;
else { /* failed to register */
result = EXIT_FAILURE;
break;
}
else ++ntuned;
}
else if (config_init->gflops < config.gflops) { /* update */
memcpy(config_init, &config, sizeof(config));
}
# if LIBXSMM_VERSION4(1, 17, 0, 0) < LIBXSMM_VERSION_NUMBER
if (active_match == i && 0 != c_dbcsr_acc_opencl_config.device.uid &&
c_dbcsr_acc_opencl_config.device.uid != key.devuid)
{
key.devuid = c_dbcsr_acc_opencl_config.device.uid;
if (active_match == i && 0 != default_uid && default_uid != key.devuid) {
key.devuid = default_uid;
config_init = (opencl_libsmm_smm_t*)libxsmm_xdispatch(&key, sizeof(key));
if (NULL == config_init && NULL != libxsmm_xregister(&key, sizeof(key), sizeof(config), &config)) {
static int info = 0;
Expand Down Expand Up @@ -532,7 +534,7 @@ int libsmm_acc_init(void) {
}
# endif
# if defined(OPENCL_KERNELS_DEVICES)
if (EXIT_SUCCESS == result && 0 != ntuned &&
if (EXIT_SUCCESS == result && 0 != ntuned && 0 == c_dbcsr_acc_opencl_config.nrank &&
(2 <= c_dbcsr_acc_opencl_config.verbosity || 0 > c_dbcsr_acc_opencl_config.verbosity))
{
fprintf(stderr, "INFO ACC/LIBSMM: PARAMS in %i set%s loaded targeting ", ntuned, 1 != ntuned ? "s" : "");
Expand Down Expand Up @@ -1084,7 +1086,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
(NULL == env_nz || '\0' == *env_nz) ? (0 != defaults ? /*default*/ 0 : config->nz) : atoi(env_nz), 0, 1);
new_config.al = LIBXSMM_CLMP(/* bug: AL=1 */
(NULL == env_al || '\0' == *env_al)
? (0 == (64 & c_dbcsr_acc_opencl_config.wa) ? (0 != defaults ? 0 : config->al) : 0)
? (0 == (32 & c_dbcsr_acc_opencl_config.wa) ? (0 != defaults ? 0 : config->al) : 0)
: atoi(env_al),
0, 1);
new_config.tb = LIBXSMM_CLMP(
Expand All @@ -1095,7 +1097,7 @@ int libsmm_acc_process(const int* host_param_stack, const int* dev_param_stack,
(NULL == env_ap || '\0' == *env_ap) ? (0 != defaults ? /*default*/ 0 : config->ap) : atoi(env_ap), 0, 1);
new_config.aa = LIBXSMM_CLMP(/* bug: AA=2 XF=1 */
(NULL == env_aa || '\0' == *env_aa) ? (0 != defaults ? default_aa : config->aa) : atoi(env_aa), 0,
(0 == (32 & c_dbcsr_acc_opencl_config.wa) || 0 == new_config.flags) ? 2 : 1);
(0 == (16 & c_dbcsr_acc_opencl_config.wa) || 0 == new_config.flags) ? 2 : 1);
new_config.ab = LIBXSMM_CLMP(
(NULL == env_ab || '\0' == *env_ab) ? (0 != defaults ? default_ab : config->ab) : atoi(env_ab), 0, 2);
new_config.ac = LIBXSMM_CLMP(
Expand Down
8 changes: 4 additions & 4 deletions src/acc/opencl/smm/params/tune_multiply_Mi250.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@ gfx90a [0x989f];3;8;5;8;30000;0;15;8;1;1;1;1;-2;0;0;0;1;0;0;0;0;0
gfx90a [0x989f];3;8;6;6;30000;0;12;8;1;8;1;-2;-2;0;0;0;1;0;2;0;0;0
gfx90a [0x989f];3;8;6;8;30000;0;15;8;1;1;1;-2;-1;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;8;7;7;30000;0;13;8;1;5;1;1;-2;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;8;7;8;30000;0;20;8;1;4;1;-2;0;0;0;0;1;1;0;2;1;0
gfx90a [0x989f];3;8;7;8;30000;0;20;8;1;6;1;-1;-2;0;0;0;1;0;0;0;0;0
gfx90a [0x989f];3;8;8;3;30000;0;15;8;1;1;1;1;-1;0;0;0;1;0;2;0;0;0
gfx90a [0x989f];3;8;8;4;30000;0;15;8;1;7;1;0;-1;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;8;8;5;30000;0;15;8;1;8;1;0;-1;0;0;0;1;0;2;2;0;0
Expand Down Expand Up @@ -212,7 +212,7 @@ gfx90a [0x989f];3;9;9;23;30000;0;15;9;1;1;1;-1;-1;0;1;1;1;1;2;2;0;0
gfx90a [0x989f];3;9;16;9;30000;0;24;9;1;8;1;0;-2;0;0;0;1;0;2;2;0;0
gfx90a [0x989f];3;9;16;16;30000;0;12;9;1;1;1;1;-2;0;0;0;1;0;2;2;0;0
gfx90a [0x989f];3;9;23;9;30000;0;26;9;1;8;1;0;-2;0;0;0;1;0;0;2;0;0
gfx90a [0x989f];3;9;23;23;30000;0;3;8;1;8;1;1;3;0;1;0;1;0;1;2;1;0
gfx90a [0x989f];3;9;23;23;30000;0;3;8;1;5;1;0;2;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;10;3;3;30000;0;10;10;1;8;1;0;-2;0;1;0;1;1;2;0;0;0
gfx90a [0x989f];3;10;3;10;30000;0;12;10;1;1;1;-1;-2;0;0;1;1;1;2;0;0;0
gfx90a [0x989f];3;10;4;4;30000;0;13;10;1;3;1;-2;-1;1;1;0;1;1;0;2;0;0
Expand Down Expand Up @@ -249,7 +249,7 @@ gfx90a [0x989f];3;11;11;3;30000;0;22;11;1;2;1;-1;1;0;1;0;1;1;2;0;1;0
gfx90a [0x989f];3;11;11;4;30000;0;20;11;1;2;1;-1;-2;0;1;1;1;1;0;0;0;0
gfx90a [0x989f];3;11;11;5;30000;0;20;11;1;5;1;-2;1;0;1;1;1;1;2;0;1;0
gfx90a [0x989f];3;11;11;6;30000;0;24;11;1;1;1;-2;0;0;0;0;1;1;2;2;0;0
gfx90a [0x989f];3;11;11;7;30000;0;9;8;1;7;1;0;1;0;1;0;1;0;1;2;0;0
gfx90a [0x989f];3;11;11;7;30000;0;9;8;1;9;1;1;1;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;11;11;11;30000;0;12;11;1;10;1;-2;-2;0;1;0;1;1;2;0;0;0
gfx90a [0x989f];3;11;11;16;30000;0;3;8;1;11;1;1;-1;1;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;11;11;23;30000;0;15;11;1;1;1;1;-1;0;1;1;1;1;2;2;0;0
Expand Down Expand Up @@ -351,7 +351,7 @@ gfx90a [0x989f];3;18;23;18;30000;0;4;18;1;4;1;1;3;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;18;23;23;30000;0;4;8;1;16;1;0;3;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;19;19;19;30000;0;40;8;1;10;1;1;3;0;0;0;1;0;1;2;0;0
gfx90a [0x989f];3;19;19;23;30000;0;40;8;1;15;1;-1;-2;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;23;23;23;30000;0;4;8;1;22;23;-1;3;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;23;23;23;30000;0;4;8;1;22;23;1;0;0;1;0;1;0;1;0;0;0
gfx90a [0x989f];3;28;28;28;30000;0;3;28;1;28;28;-2;2;0;0;0;1;0;1;0;0;0
gfx90a [0x989f];3;32;32;32;30000;0;25;32;1;20;1;-2;0;0;1;0;1;0;2;0;0;0
gfx90a [0x989f];3;35;17;17;30000;0;15;35;1;29;1;1;0;0;1;0;1;0;2;1;0;0
Expand Down
Loading
Loading